//#pragma once
//// CUDA Runtime
//#include <cuda_runtime.h>
//#include <device_functions.h>
//#include <device_launch_parameters.h>
//// Utilities and system includes
//#include <helper_cuda.h>
//#include <helper_functions.h>
//#include <thrust/device_vector.h>
//#include <thrust/scan.h>
//#include <iostream>
//#include <math.h>
//#include "XModel.h"
//#include "math_functions.hpp"
//#include "cuda_runtime_api.h"
//
//using namespace cudacp;
//
//#ifndef MIN
//#define MIN(x,y) ((x < y) ? x : y)
//#endif
//// num_threads
//static const int num_threads = 128;
//static const int U32_SIZE = sizeof(u32); ///<4
//static const int U32_BIT = U32_SIZE * 8;	///<32
//static const int U32_POS = 5;
//static const int U32_MOD_MASK = 31;
//
//struct int_predicate
//{
//	__host__ __device__	bool operator()(const int x)
//	{
//		return x > 0;
//	}
//};
//
//// һbitDom[x]ĳ
//__constant__ int D_BITDOM_INTSIZE;
//// bitDomĳ
//
//__constant__ int D_BITDOMS_INTSIZE;
//
//// bit֧֣ uint2 bitSup[c][a][idx].x = bitSup[c,x,a,idx]
//// bit֧֣ uint2 bitSup[c][a][idx].y = bitSup[c,y,a,idx]
////__device__ uint2*** d_bitSup;
////__host__ uint2*** h_bitSup;
////__device__ u32** d_bitDom;
////__host__ u32** h_bitDom;
////__device__ u32**
//__constant__ int D_NUM_BD_BLOCK;
//__constant__ int D_NUM_CS_SIZE_BLOCKS;
////////////////////////////////////////////////////////////////////////////
////	һЩGPU
////////////////////////////////////////////////////////////////////////////
//__device__ __managed__ int NUM_BD_BLOCK;
//__device__ __managed__ int NUM_CS_SIZE_BLOCKS;
//// һint
//__device__ __managed__ int BITDOM_INTSIZE;
//// ϵint 
//__device__ __managed__ int BITDOMS_INTSIZE;
//// һԼbitsupint 
//__device__ __managed__ int BITSUP_INTSIZE;
//// Լϵbitsupint
//__device__ __managed__ int BITSUPS_INTSIZE;
////ܳ
//__device__ __managed__ int BITSUBDOMS_INTSIZE;
////
//__device__ __managed__ int VS_SIZE;
////Լ
//__device__ __managed__ int CS_SIZE;
////ԼѹBLOCK
//__device__ __managed__ int MCC_BLOCK;
////////////////////////////////////////////////////////////////////////////
////	һЩGPU
////////////////////////////////////////////////////////////////////////////
//__device__ __managed__ int M_Qsize;
//
////////////////////////////////////////////////////////////////////////////
////  GPUԼ¼Ϣɸ
////////////////////////////////////////////////////////////////////////////
////	ÿintС
//__device__ __managed__ int *vars_size;
//// 洢Լscopeint3scope.x: x.id; scope.y: y.id; scope.z: c.id
//__device__ __managed__ int3* scope;
//// dom
//__device__ __managed__ int MAX_DOM_SIZE;
//// subCon
//__device__ __managed__ int SUBCON_SIZE;
//
//
////__device__ __managed__ int BITDOM_SIZE;
////__device__ __managed__ int
//// ݽṹʹUM
//// ʾԼ
//__device__ __managed__ u32* bitDom;
//// ʾԼ޸
//__device__ __managed__ uint2* bitSup;
//////ƶУ洢Լid
////__device__ __managed__ int *mainCon;
//////ݽṹ
////ʾԼ
//__device__ __managed__ u32* bitSubDom;
//////ƶУ洢Լid subCon.x: variablesubCon.y: valuesubCon.z: c.id
////__device__ __managed__ ushort3* subCon;
////ǷɾʼȫΪ1
//__device__ __managed__ int* M_VarPre;
////ԼǷ飬ʼȫΪ1
//__device__ __managed__ int* M_ConPre;
////Լ(ѹ)
//__device__ __managed__ int* M_ConEvt;
////Լ
//__device__ __managed__ int* M_Con;
//
////ǷɾʼȫΪ1
//__device__ __managed__ int* S_VarPre;
////¼
//__device__ __managed__ uint3* S_Var;
////ԼǷ飬ʼȫΪ1
//__device__ __managed__ int* S_ConPre;
////Լ(ѹ)
//__device__ __managed__ int3* S_ConEvt;
////Լ
//__device__ __managed__ int3* S_Con;
////
//int* MCC_BlocksCount;
//int* MCC_BlocksOffset;
//
//thrust::device_vector<int> MCC_BCount;
//thrust::device_vector<int> MCC_BOffset;
////
////__device__ __managed__ ushort4* subVar;
//////ⷢĶıidʼȫΪ0
////__device__ __managed__ unsigned short* subEvtVar;
//////ⷢĶԼidʼȫΪ1
////__device__ __managed__ int* subEvtCon;
//
//static const u32 U32_MASK1[32] = {
//	0x80000000, 0x40000000, 0x20000000, 0x10000000,
//	0x08000000, 0x04000000, 0x02000000, 0x01000000,
//	0x00800000, 0x00400000, 0x00200000, 0x00100000,
//	0x00080000, 0x00040000, 0x00020000, 0x00010000,
//	0x00008000, 0x00004000, 0x00002000, 0x00001000,
//	0x00000800, 0x00000400, 0x00000200, 0x00000100,
//	0x00000080, 0x00000040, 0x00000020, 0x00000010,
//	0x00000008, 0x00000004, 0x00000002, 0x00000001,
//};
//
//static const u32 U32_MASK0[32] = {
//	0x7FFFFFFF, 0xBFFFFFFF, 0xDFFFFFFF, 0xEFFFFFFF,
//	0xF7FFFFFF, 0xFBFFFFFF, 0xFDFFFFFF, 0xFEFFFFFF,
//	0xFF7FFFFF, 0xFFBFFFFF, 0xFFDFFFFF, 0xFFEFFFFF,
//	0xFFF7FFFF, 0xFFFBFFFF, 0xFFFDFFFF, 0xFFFEFFFF,
//	0xFFFF7FFF, 0xFFFFBFFF, 0xFFFFDFFF, 0xFFFFEFFF,
//	0xFFFFF7FF, 0xFFFFFBFF, 0xFFFFFDFF, 0xFFFFFEFF,
//	0xFFFFFF7F, 0xFFFFFFBF, 0xFFFFFFDF, 0xFFFFFFEF,
//	0xFFFFFFF7, 0xFFFFFFFB, 0xFFFFFFFD, 0xFFFFFFFE,
//};
//
////__forceinline__ int  GetBitDomIndex(int var_id)
////{
////	return var_id * BITDOM_INTSIZE;
////}
//
//// xindexbitDomλ
//#define GetBitDomIndex(x, i) (x * BITDOM_INTSIZE + i)
//// ֵֵĸȡbitʾƫ
//#define GetOffSet(x)(U32_BIT - (x & U32_MOD_MASK))
//
//#define GetBitSubDomStartIndex(x,a)((x * MAX_DOM_SIZE + a) * BITDOMS_INTSIZE)
//#define GetBitSubDomIndex(x, a, y, i)(GetBitSubDomStartIndex(x,a) + GetBitDomIndex(y, i))
//
//__device__ bool IsGtZero(int x)
//{
//	return x > 0;
//}
//
//__inline__ __device__ __host__ int GetTopNum(int num_elements, int num_threads)
//{
//	return (num_elements + (num_threads - 1)) / num_threads;
//}
////************************************
//// Method:    intsizeof
//// FullName:  intsizeof
//// Access:    public 
//// Returns:   int
//// Qualifier: ȡbitʾint
//// Parameter: const int x
////************************************
//inline int intsizeof(const int x)
//{
//	return (int)ceil((float)x / U32_BIT);
//}
//
//__device__ __inline__ int pow2i(int e)
//{
//	return 1 << e;
//}
////__global__ void enforceACMain(u32* bitDom, u32* bitSup, u32* M_Con, u32* M_ConEvt, u32* M_ConPre, u32* M_VarPre)
////{
////
////}
//
////ͨѸıԼ
//__global__ void GenConPre(int *VarPre, int* BlocksCount, int3* scp, int len)
//{
//	const int idx = blockDim.x*blockIdx.x + threadIdx.x;
//	if (idx < len)
//	{
//		int3 sp = scp[idx];
//		int pred;
//		if (VarPre[sp.x] == 1 || VarPre[sp.y] == 1)
//			pred = 1;
//		else
//			pred = 0;
//
//		int BC = __syncthreads_count(pred);
//
//		if (threadIdx.x == 0)
//		{
//			BlocksCount[threadIdx.x] = BC;
//		}
//	}
//}
//
//__global__ void CompactQ(int *VarPre, int* ConEvt, int* BOffset, int3* scp, int len)
//{
//	int idx = threadIdx.x + blockIdx.x*blockDim.x;
//	extern __shared__ int warpTotals[];
//	//һ߳̿128߳
//	//һ4߳
//	if (idx < len)
//	{
//		int3 sp = scp[idx];
//		int pred;
//		//ж
//		if (VarPre[sp.x] == 1 || VarPre[sp.y] == 1)
//			pred = 1;
//		else
//			pred = 0;
//
//		//warp index
//		//߳
//		int w_i = threadIdx.x / warpSize;
//		//thread index within a warp
//		//߳߳
//		int w_l = idx % warpSize;
//		//thread mask (ERROR IN THE PAPERminus one is required)
//		//߳
//		//INT_MAX = 1111 1111 1111 1111 1111 1111 1111 1111 
//		//߳id=032-0-1 = 31λ Ҳʣ1λ
//		//߳id=532-5-1 = 26λ Ҳʣ6λ
//		//߳id=3132-31-1 = 0λ Ҳʣ32λ
//		//߳threid|  31~~~~~~0
//		//ballotӦλ|   1......1
//		int t_m = INT_MAX >> (warpSize - w_l - 1);
//		//balres = number whose ith bit is one if the ith's thread pred is true masked up to the current index in warp
//		//߳ھֲpred = 1밴λ뵫˵߳idļ¼ֻǰж
//		int b = __ballot(pred) & t_m;
//		//popc count the number of bit one. simply count the number predicated true BEFORE MY INDEX
//		//ֻ㵱ǰ߳ӦǰNλ֮
//		//Ϊ߳ɨ
//		int t_u = __popc(b);
//
//		//ÿ߳һ߳д빲ڴ棬ӦidΪ߳ID߳IDӻأ
//		//͵ֵд빲ڴ棬͵ֵûб
//		//warpTotalsΪ4
//		if (w_l == warpSize - 1)
//			warpTotals[w_i] = t_u + pred;
//
//		__syncthreads();
//
//		//߳idΪ0߳߳idblockDim.x = 128w_l < 128/32 = 4
//		//߳̿ڵһ߳ǰ4߳w_l < ߳4ÿ߳һ߳
//		if (w_i == 0 && w_l < blockDim.x / warpSize)
//		{
//			int w_i_u = 0;
//			for (int j = 0; j <= 5; ++j)
//			{
//				//# of the ones in the j'th digit of the warp offsets
//				//0->5 6λã
//				//000 001
//				//000 010
//				//000 100
//				//001 000
//				//010 000
//				//100 000
//				int b_j = __ballot(warpTotals[w_l] & pow2i(j));
//				w_i_u += (__popc(b_j & t_m)) << j;
//				//printf("indice %i t_m=%i,j=%i,b_j=%i,w_i_u=%i\n",w_l,t_m,j,b_j,w_i_u);
//			}
//			warpTotals[w_l] = w_i_u;
//		}
//		__syncthreads();
//
//		if (pred)
//			ConEvt[t_u + warpTotals[w_i] + BOffset[blockIdx.x]] = scp[idx].z;
//
//	}
//}
//
//void CompactQueueMain()
//{
//	//Լ
//	//P1
//	GenConPre << <MCC_BLOCK, num_threads >> >(M_VarPre, MCC_BlocksCount, scope, CS_SIZE);
//	cudaDeviceSynchronize();
//	//P2
//	thrust::exclusive_scan(MCC_BCount.begin(), MCC_BCount.end(), MCC_BOffset.begin());
//	cudaDeviceSynchronize();
//	//P3
//	//ÿԼһ߳̽йԼ,ڴС = һ߳ĸ,װ߳
//	CompactQ << <MCC_BLOCK, num_threads, sizeof(int)*(num_threads / warpSize) >> >(M_VarPre, M_Con, MCC_BlocksOffset, scope, CS_SIZE);
//}
//
//#define GetBitSupIndexByINTPrstn(cid,x_val,y_val) (cid * BITSUP_INTSIZE + x_val * BITDOM_INTSIZE + y_val)
//
//__inline__ __device__ __host__ int2 GetBitSupIndexByTuple(int cid, int2 t)
//{
//	return make_int2(cid * BITSUP_INTSIZE + t.x * BITDOM_INTSIZE + (t.y >> U32_POS), cid * BITSUP_INTSIZE + t.y * BITDOM_INTSIZE + (t.x >> U32_POS));
//}
//
//__inline__ __device__ __host__ int GetBitSupIndexById(int cid)
//{
//	return cid * BITSUP_INTSIZE;
//}
//
//void DelGPUModel();
//
//void BuildBitModel(XModel *xm)
//{
//#pragma region 㳣
//	BITDOM_INTSIZE = intsizeof(xm->feature.max_dom_size);
//	MAX_DOM_SIZE = xm->feature.max_dom_size;
//	VS_SIZE = xm->feature.vs_size;
//	CS_SIZE = xm->feature.cs_size;
//	BITDOMS_INTSIZE = BITDOM_INTSIZE * VS_SIZE;
//	BITSUP_INTSIZE = MAX_DOM_SIZE * BITDOM_INTSIZE;
//	BITSUPS_INTSIZE = BITSUP_INTSIZE * CS_SIZE;
//	BITSUBDOMS_INTSIZE = VS_SIZE * MAX_DOM_SIZE * BITDOMS_INTSIZE;
//	SUBCON_SIZE = VS_SIZE * MAX_DOM_SIZE * CS_SIZE;
//#pragma endregion 㳣
//#pragma region ԼϢ
//	//cudaMallocManaged(&vars_size, sizeof(int) * VS_SIZE);
//	//// ʼС
//	//for (int i = 0; i < xm->feature.vs_size; ++i)
//	//{
//	//	XVar* v = xm->vars[i];
//	//	XDom* d = xm->doms[v->dom_id];
//	//	vars_size[i] = d->size;
//	//}
//
//	// ʼscope
//	cudaMallocManaged(&scope, sizeof(int3) * CS_SIZE);
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		XCon *c = xm->cons[i];
//		scope[i].x = c->scope[0];
//		scope[i].y = c->scope[1];
//		scope[i].z = c->id;
//	}
//
//	////ʾ
//	//for (int i = 0; i < CS_SIZE; ++i)
//	//{
//	//	printf("scope[%d] = {%d, %d}\n", scope[i].z, scope[i].x, scope[i].y);
//	//}
//#pragma endregion ԼϢ
//#pragma region bitDom
//	cudaMallocManaged(&bitDom, sizeof(u32) * BITDOMS_INTSIZE);
//	cudaMallocManaged(&M_VarPre, sizeof(int) * VS_SIZE);
//
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		XVar* v = xm->vars[i];
//		XDom* d = xm->doms[v->dom_id];
//		const int dom_size = d->size;
//		// ǰʵINT
//		const int dom_int_size = intsizeof(dom_size);
//
//		for (int j = 0; j < BITDOM_INTSIZE; ++j)
//		{
//			const int idx = GetBitDomIndex(i, j);
//			//printf("idx = %d\n", idx);
//			// 
//			if (j < dom_int_size - 1)
//				bitDom[idx] = UINT32_MAX;
//			else if (j == dom_int_size - 1)
//				bitDom[idx] = UINT32_MAX << GetOffSet(dom_size);
//			else
//				bitDom[idx] = 0;
//		}
//
//		M_VarPre[i] = 1;
//	}
//
//	//for (int i = 0; i < VS_SIZE; ++i)
//	//{
//	//	for (int j = 0; j < BITDOM_INTSIZE; ++j)
//	//	{
//	//		int idx = GetBitDomIndex(i, j);
//	//		printf("var = %d, j = %d, idx = %d, bitDom = %x, pre= %x\n", i, j, idx, bitDom[idx], M_VarPre[i]);
//	//	}
//	//}
//#pragma endregion bitDom
//#pragma region bitSubDom
//	cudaMallocManaged(&bitSubDom, sizeof(u32)*BITDOMS_INTSIZE*VS_SIZE*MAX_DOM_SIZE);
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			const int start_idx = GetBitSubDomStartIndex(i, j);
//			for (int k = 0; k < BITSUBDOMS_INTSIZE; ++k)
//				bitSubDom[start_idx + k] = bitDom[k];
//			//(i,j)bitDom ĵ
//			//ȡi,j,iʼַ
//			const int ijistart = start_idx + i*BITDOM_INTSIZE;
//			for (int k = 0; k < BITDOM_INTSIZE; ++k)
//				// jKķΧ:j/32,j%32λΪ1
//				if (k == j >> U32_POS)
//					bitSubDom[ijistart + k] = U32_MASK1[j&U32_MOD_MASK];
//			//λΪ0
//				else
//					bitSubDom[ijistart + k] = 0;
//		}
//	}
//
//	//for (int i = 0; i < VS_SIZE; ++i)
//	//{
//	//	for (int j = 0; j < MAX_DOM_SIZE; ++j)
//	//	{
//	//		printf("sub problem:(%d, %d): ", i, j);
//	//		const int start_idx = GetBitSubDomStartIndex(i, j);
//	//		for (int k = 0; k < BITDOMS_INTSIZE; ++k)
//	//		{
//	//			printf("%x ", bitSubDom[start_idx + k]);
//	//		}
//	//		printf("\n");
//	//	}
//	//}
//#pragma endregion bitSubDom
//#pragma region bitSup
//	cudaMallocManaged(&bitSup, sizeof(uint2) * BITSUPS_INTSIZE);
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		XCon* c = xm->cons[i];
//		XRel* r = xm->rels[c->rel_id];
//		XVar* v[2] = { xm->vars[c->scope[0]], xm->vars[c->scope[1]] };
//		XDom* d[2] = { xm->doms[v[0]->dom_id], xm->doms[v[1]->dom_id] };
//
//		//ʼλ
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			for (int k = 0; k < BITDOM_INTSIZE; ++k)
//			{
//				const int idx = GetBitSupIndexByINTPrstn(c->id, j, k);
//				if (j < d[0]->size && (k < (d[1]->size >> U32_POS)))
//				{
//					//֧ȡ0x0000..., ͻȡ0xFFF...
//					bitSup[idx].x = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].y = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//				}
//				else if (k == (d[1]->size >> U32_POS))
//				{
//					bitSup[idx].x = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].y = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].x <<= U32_BIT - (d[1]->size & U32_MOD_MASK);
//					bitSup[idx].y <<= U32_BIT - (d[1]->size & U32_MOD_MASK);
//				}
//				else
//				{
//					bitSup[idx].x = 0;
//					bitSup[idx].y = 0;
//				}
//			}
//		}
//		//λֵ
//		for (int j = 0; j < r->size; ++j)
//		{
//			const int2 t = make_int2(r->tuples[j][0], r->tuples[j][1]);
//			//printf("c_id= %d, %d, %d\n", c->id, t.x, t.y);
//			const int2 idx = GetBitSupIndexByTuple(c->id, t);
//			//printf("idx = %d, %d\n", idx.x, idx.y);
//			if (r->sem == SEM_SUPPORT)
//			{
//				bitSup[idx.x].x |= U32_MASK1[t.y & U32_MOD_MASK];
//				bitSup[idx.y].y |= U32_MASK1[t.x & U32_MOD_MASK];
//			}
//			else
//			{
//				bitSup[idx.x].x &= U32_MASK0[t.y & U32_MOD_MASK];
//				bitSup[idx.y].y &= U32_MASK0[t.x & U32_MOD_MASK];
//			}
//		}
//		//// ʼλ
//		//for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		//{
//		//	printf("c_id = %d, j = %d: ", i, j);
//		//	for (int k = 0; k < BITDOM_INTSIZE; ++k)
//		//	{
//		//		const int idx = GetBitSupIndexByINTPrstn(c->id, j, k);
//		//		printf("%x, %x", bitSup[idx].x, bitSup[idx].y);
//		//	}
//		//	printf("\n");
//		//}
//	}
//#pragma endregion bitSup
//#pragma region Լ
//	cudaMallocManaged(&M_Con, sizeof(int)*CS_SIZE);
//	cudaMallocManaged(&M_ConEvt, sizeof(int) * CS_SIZE);
//	cudaMallocManaged(&M_ConPre, sizeof(int)*CS_SIZE);
//
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		M_Con[i] = i;
//		M_ConEvt[i] = i;
//		M_ConPre[i] = 1;
//		//printf("i = %d , M_Con = %d, M_ConEvt = %d, M_ConPre = %d\n", i, M_Con[i], M_ConEvt[i], M_ConPre[i]);
//	}
//#pragma endregion Լ
//#pragma region Լ
//	cudaMallocManaged(&S_ConPre, sizeof(int)*SUBCON_SIZE);
//	cudaMallocManaged(&S_ConEvt, sizeof(int3)*SUBCON_SIZE);
//	cudaMallocManaged(&S_Con, sizeof(int3)*SUBCON_SIZE);
//	cudaMallocManaged(&S_VarPre, sizeof(int)*VS_SIZE*MAX_DOM_SIZE*VS_SIZE);
//	cudaMallocManaged(&S_Var, sizeof(int3)*VS_SIZE*MAX_DOM_SIZE*VS_SIZE);
//
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		const XVar* v = xm->vars[i];
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			for (int k = 0; k < CS_SIZE; ++k)
//			{
//				// (i, j) kΪԼid
//				const int idx = (i*MAX_DOM_SIZE + j)*CS_SIZE + k;
//				//i*xm->feature.max_dom_size*xm->feature.cs_size + j*xm->feature.cs_size + k;
//
//				S_Con[idx].x = i;
//				S_Con[idx].y = j;
//				S_Con[idx].z = k;
//
//				S_ConEvt[idx].x = i;
//				S_ConEvt[idx].y = j;
//				S_ConEvt[idx].z = k;
//
//				S_ConPre[idx] = 1;
//				//printf("S_Con = (%d, %d, %d), S_ConEvt = (%d, %d, %d), pre = %d\n", S_Con[idx].x, S_Con[idx].y, S_Con[idx].z, S_ConEvt[idx].x, S_ConEvt[idx].y, S_ConEvt[idx].z, S_ConPre[idx]);
//			}
//
//			for (int k = 0; k < VS_SIZE; ++k)
//			{
//				//(i, j) kΪid
//				const int idx = (i*MAX_DOM_SIZE + j)*VS_SIZE + k;
//				S_Var[idx].x = i;
//				S_Var[idx].y = j;
//				S_Var[idx].z = k;
//
//				S_VarPre[idx] = 1;
//
//				//printf("S_Var = (%d, %d, %d), S_VarPre = %d\n", S_Var[idx].x, S_Var[idx].y, S_Var[idx].z, S_VarPre[idx]);
//			}
//		}
//	}
//#pragma endregion Լ
//
//#pragma region й
//	//ѹBLOCK
//	MCC_BLOCK = GetTopNum(CS_SIZE, num_threads);
//	MCC_BCount.resize(MCC_BLOCK, 0);
//	MCC_BOffset.resize(MCC_BLOCK, 0);
//	MCC_BlocksCount = thrust::raw_pointer_cast(MCC_BCount.data());
//	MCC_BlocksOffset = thrust::raw_pointer_cast(MCC_BOffset.data());
//#pragma endregion
//
//}
//
//__global__ void ConCheckMain(int* ConEvt, int* btSp, int2* scp)
//{
//	const int c_id = blockIdx.x;
//	//ȡԼbitSupĿʼ
//	const int start_idx = GetBitSupIndexById(c_id);
//	const int2 sp = scp[c_id];
//	extern __shared__ int2[];
//}
//
//void ConstraintCheckMain()
//{
//	//num_threadsÿɱ
//	ConCheckMain << <CS_SIZE, num_threads >> >();
//}
//
//float SACGPU()
//{
//	//1. ִAC
//	//1.1. ѹ
//	CompactQueueMain();
//	//1.2. Լ
//	ConstraintCheckMain();
//}
//
//void DelGPUModel()
//{
//	cudaFree(scope);
//	cudaFree(bitDom);
//	cudaFree(M_VarPre);
//	cudaFree(bitSubDom);
//	cudaFree(bitSup);
//	cudaFree(M_Con);
//	cudaFree(M_ConEvt);
//	cudaFree(M_ConPre);
//	cudaFree(S_ConPre);
//	cudaFree(S_ConEvt);
//	cudaFree(S_Con);
//	cudaFree(S_Var);
//	cudaFree(S_VarPre);
//}
//
//

	.file	"tmpxft_002c4b69_00000000-6_cudasac.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2033:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
//#pragma once
//// CUDA Runtime
//#include <cuda_runtime.h>
//#include <device_functions.h>
//#include <device_launch_parameters.h>
//// Utilities and system includes
//#include <helper_cuda.h>
//#include <helper_functions.h>
//#include <thrust/device_vector.h>
//#include <thrust/scan.h>
//#include <iostream>
//#include <math.h>
//#include "XModel.h"
//#include "math_functions.hpp"
//#include "cuda_runtime_api.h"
//
//using namespace cudacp;
//
//#ifndef MIN
//#define MIN(x,y) ((x < y) ? x : y)
//#endif
//// num_threads
//static const int num_threads = 128;
//static const int U32_SIZE = sizeof(u32); ///<4
//static const int U32_BIT = U32_SIZE * 8;	///<32
//static const int U32_POS = 5;
//static const int U32_MOD_MASK = 31;
//
//struct int_predicate
//{
//	__host__ __device__	bool operator()(const int x)
//	{
//		return x > 0;
//	}
//};
//
//// һbitDom[x]ĳ
//__constant__ int D_BITDOM_INTSIZE;
//// bitDomĳ
//
//__constant__ int D_BITDOMS_INTSIZE;
//
//// bit֧֣ uint2 bitSup[c][a][idx].x = bitSup[c,x,a,idx]
//// bit֧֣ uint2 bitSup[c][a][idx].y = bitSup[c,y,a,idx]
////__device__ uint2*** d_bitSup;
////__host__ uint2*** h_bitSup;
////__device__ u32** d_bitDom;
////__host__ u32** h_bitDom;
////__device__ u32**
//__constant__ int D_NUM_BD_BLOCK;
//__constant__ int D_NUM_CS_SIZE_BLOCKS;
////////////////////////////////////////////////////////////////////////////
////	һЩGPU
////////////////////////////////////////////////////////////////////////////
//__device__ __managed__ int NUM_BD_BLOCK;
//__device__ __managed__ int NUM_CS_SIZE_BLOCKS;
//// һint
//__device__ __managed__ int BITDOM_INTSIZE;
//// ϵint 
//__device__ __managed__ int BITDOMS_INTSIZE;
//// һԼbitsupint 
//__device__ __managed__ int BITSUP_INTSIZE;
//// Լϵbitsupint
//__device__ __managed__ int BITSUPS_INTSIZE;
////ܳ
//__device__ __managed__ int BITSUBDOMS_INTSIZE;
////
//__device__ __managed__ int VS_SIZE;
////Լ
//__device__ __managed__ int CS_SIZE;
////ԼѹBLOCK
//__device__ __managed__ int MCC_BLOCK;
////////////////////////////////////////////////////////////////////////////
////	һЩGPU
////////////////////////////////////////////////////////////////////////////
//__device__ __managed__ int M_Qsize;
//
////////////////////////////////////////////////////////////////////////////
////  GPUԼ¼Ϣɸ
////////////////////////////////////////////////////////////////////////////
////	ÿintС
//__device__ __managed__ int *vars_size;
//// 洢Լscopeint3scope.x: x.id; scope.y: y.id; scope.z: c.id
//__device__ __managed__ int3* scope;
//// dom
//__device__ __managed__ int MAX_DOM_SIZE;
//// subCon
//__device__ __managed__ int SUBCON_SIZE;
//
//
////__device__ __managed__ int BITDOM_SIZE;
////__device__ __managed__ int
//// ݽṹʹUM
//// ʾԼ
//__device__ __managed__ u32* bitDom;
//// ʾԼ޸
//__device__ __managed__ uint2* bitSup;
//////ƶУ洢Լid
////__device__ __managed__ int *mainCon;
//////ݽṹ
////ʾԼ
//__device__ __managed__ u32* bitSubDom;
//////ƶУ洢Լid subCon.x: variablesubCon.y: valuesubCon.z: c.id
////__device__ __managed__ ushort3* subCon;
////ǷɾʼȫΪ1
//__device__ __managed__ int* M_VarPre;
////ԼǷ飬ʼȫΪ1
//__device__ __managed__ int* M_ConPre;
////Լ(ѹ)
//__device__ __managed__ int* M_ConEvt;
////Լ
//__device__ __managed__ int* M_Con;
//
////ǷɾʼȫΪ1
//__device__ __managed__ int* S_VarPre;
////¼
//__device__ __managed__ uint3* S_Var;
////ԼǷ飬ʼȫΪ1
//__device__ __managed__ int* S_ConPre;
////Լ(ѹ)
//__device__ __managed__ int3* S_ConEvt;
////Լ
//__device__ __managed__ int3* S_Con;
////
//int* MCC_BlocksCount;
//int* MCC_BlocksOffset;
//
//thrust::device_vector<int> MCC_BCount;
//thrust::device_vector<int> MCC_BOffset;
////
////__device__ __managed__ ushort4* subVar;
//////ⷢĶıidʼȫΪ0
////__device__ __managed__ unsigned short* subEvtVar;
//////ⷢĶԼidʼȫΪ1
////__device__ __managed__ int* subEvtCon;
//
//static const u32 U32_MASK1[32] = {
//	0x80000000, 0x40000000, 0x20000000, 0x10000000,
//	0x08000000, 0x04000000, 0x02000000, 0x01000000,
//	0x00800000, 0x00400000, 0x00200000, 0x00100000,
//	0x00080000, 0x00040000, 0x00020000, 0x00010000,
//	0x00008000, 0x00004000, 0x00002000, 0x00001000,
//	0x00000800, 0x00000400, 0x00000200, 0x00000100,
//	0x00000080, 0x00000040, 0x00000020, 0x00000010,
//	0x00000008, 0x00000004, 0x00000002, 0x00000001,
//};
//
//static const u32 U32_MASK0[32] = {
//	0x7FFFFFFF, 0xBFFFFFFF, 0xDFFFFFFF, 0xEFFFFFFF,
//	0xF7FFFFFF, 0xFBFFFFFF, 0xFDFFFFFF, 0xFEFFFFFF,
//	0xFF7FFFFF, 0xFFBFFFFF, 0xFFDFFFFF, 0xFFEFFFFF,
//	0xFFF7FFFF, 0xFFFBFFFF, 0xFFFDFFFF, 0xFFFEFFFF,
//	0xFFFF7FFF, 0xFFFFBFFF, 0xFFFFDFFF, 0xFFFFEFFF,
//	0xFFFFF7FF, 0xFFFFFBFF, 0xFFFFFDFF, 0xFFFFFEFF,
//	0xFFFFFF7F, 0xFFFFFFBF, 0xFFFFFFDF, 0xFFFFFFEF,
//	0xFFFFFFF7, 0xFFFFFFFB, 0xFFFFFFFD, 0xFFFFFFFE,
//};
//
////__forceinline__ int  GetBitDomIndex(int var_id)
////{
////	return var_id * BITDOM_INTSIZE;
////}
//
//// xindexbitDomλ
//#define GetBitDomIndex(x, i) (x * BITDOM_INTSIZE + i)
//// ֵֵĸȡbitʾƫ
//#define GetOffSet(x)(U32_BIT - (x & U32_MOD_MASK))
//
//#define GetBitSubDomStartIndex(x,a)((x * MAX_DOM_SIZE + a) * BITDOMS_INTSIZE)
//#define GetBitSubDomIndex(x, a, y, i)(GetBitSubDomStartIndex(x,a) + GetBitDomIndex(y, i))
//
//__device__ bool IsGtZero(int x)
//{
//	return x > 0;
//}
//
//__inline__ __device__ __host__ int GetTopNum(int num_elements, int num_threads)
//{
//	return (num_elements + (num_threads - 1)) / num_threads;
//}
////************************************
//// Method:    intsizeof
//// FullName:  intsizeof
//// Access:    public 
//// Returns:   int
//// Qualifier: ȡbitʾint
//// Parameter: const int x
////************************************
//inline int intsizeof(const int x)
//{
//	return (int)ceil((float)x / U32_BIT);
//}
//
//__device__ __inline__ int pow2i(int e)
//{
//	return 1 << e;
//}
////__global__ void enforceACMain(u32* bitDom, u32* bitSup, u32* M_Con, u32* M_ConEvt, u32* M_ConPre, u32* M_VarPre)
////{
////
////}
//
////ͨѸıԼ
//__global__ void GenConPre(int *VarPre, int* BlocksCount, int3* scp, int len)
//{
//	const int idx = blockDim.x*blockIdx.x + threadIdx.x;
//	if (idx < len)
//	{
//		int3 sp = scp[idx];
//		int pred;
//		if (VarPre[sp.x] == 1 || VarPre[sp.y] == 1)
//			pred = 1;
//		else
//			pred = 0;
//
//		int BC = __syncthreads_count(pred);
//
//		if (threadIdx.x == 0)
//		{
//			BlocksCount[threadIdx.x] = BC;
//		}
//	}
//}
//
//__global__ void CompactQ(int *VarPre, int* ConEvt, int* BOffset, int3* scp, int len)
//{
//	int idx = threadIdx.x + blockIdx.x*blockDim.x;
//	extern __shared__ int warpTotals[];
//	//һ߳̿128߳
//	//һ4߳
//	if (idx < len)
//	{
//		int3 sp = scp[idx];
//		int pred;
//		//ж
//		if (VarPre[sp.x] == 1 || VarPre[sp.y] == 1)
//			pred = 1;
//		else
//			pred = 0;
//
//		//warp index
//		//߳
//		int w_i = threadIdx.x / warpSize;
//		//thread index within a warp
//		//߳߳
//		int w_l = idx % warpSize;
//		//thread mask (ERROR IN THE PAPERminus one is required)
//		//߳
//		//INT_MAX = 1111 1111 1111 1111 1111 1111 1111 1111 
//		//߳id=032-0-1 = 31λ Ҳʣ1λ
//		//߳id=532-5-1 = 26λ Ҳʣ6λ
//		//߳id=3132-31-1 = 0λ Ҳʣ32λ
//		//߳threid|  31~~~~~~0
//		//ballotӦλ|   1......1
//		int t_m = INT_MAX >> (warpSize - w_l - 1);
//		//balres = number whose ith bit is one if the ith's thread pred is true masked up to the current index in warp
//		//߳ھֲpred = 1밴λ뵫˵߳idļ¼ֻǰж
//		int b = __ballot(pred) & t_m;
//		//popc count the number of bit one. simply count the number predicated true BEFORE MY INDEX
//		//ֻ㵱ǰ߳ӦǰNλ֮
//		//Ϊ߳ɨ
//		int t_u = __popc(b);
//
//		//ÿ߳һ߳д빲ڴ棬ӦidΪ߳ID߳IDӻأ
//		//͵ֵд빲ڴ棬͵ֵûб
//		//warpTotalsΪ4
//		if (w_l == warpSize - 1)
//			warpTotals[w_i] = t_u + pred;
//
//		__syncthreads();
//
//		//߳idΪ0߳߳idblockDim.x = 128w_l < 128/32 = 4
//		//߳̿ڵһ߳ǰ4߳w_l < ߳4ÿ߳һ߳
//		if (w_i == 0 && w_l < blockDim.x / warpSize)
//		{
//			int w_i_u = 0;
//			for (int j = 0; j <= 5; ++j)
//			{
//				//# of the ones in the j'th digit of the warp offsets
//				//0->5 6λã
//				//000 001
//				//000 010
//				//000 100
//				//001 000
//				//010 000
//				//100 000
//				int b_j = __ballot(warpTotals[w_l] & pow2i(j));
//				w_i_u += (__popc(b_j & t_m)) << j;
//				//printf("indice %i t_m=%i,j=%i,b_j=%i,w_i_u=%i\n",w_l,t_m,j,b_j,w_i_u);
//			}
//			warpTotals[w_l] = w_i_u;
//		}
//		__syncthreads();
//
//		if (pred)
//			ConEvt[t_u + warpTotals[w_i] + BOffset[blockIdx.x]] = scp[idx].z;
//
//	}
//}
//
//void CompactQueueMain()
//{
//	//Լ
//	//P1
//	GenConPre << <MCC_BLOCK, num_threads >> >(M_VarPre, MCC_BlocksCount, scope, CS_SIZE);
//	cudaDeviceSynchronize();
//	//P2
//	thrust::exclusive_scan(MCC_BCount.begin(), MCC_BCount.end(), MCC_BOffset.begin());
//	cudaDeviceSynchronize();
//	//P3
//	//ÿԼһ߳̽йԼ,ڴС = һ߳ĸ,װ߳
//	CompactQ << <MCC_BLOCK, num_threads, sizeof(int)*(num_threads / warpSize) >> >(M_VarPre, M_Con, MCC_BlocksOffset, scope, CS_SIZE);
//}
//
//#define GetBitSupIndexByINTPrstn(cid,x_val,y_val) (cid * BITSUP_INTSIZE + x_val * BITDOM_INTSIZE + y_val)
//
//__inline__ __device__ __host__ int2 GetBitSupIndexByTuple(int cid, int2 t)
//{
//	return make_int2(cid * BITSUP_INTSIZE + t.x * BITDOM_INTSIZE + (t.y >> U32_POS), cid * BITSUP_INTSIZE + t.y * BITDOM_INTSIZE + (t.x >> U32_POS));
//}
//
//__inline__ __device__ __host__ int GetBitSupIndexById(int cid)
//{
//	return cid * BITSUP_INTSIZE;
//}
//
//void DelGPUModel();
//
//void BuildBitModel(XModel *xm)
//{
//#pragma region 㳣
//	BITDOM_INTSIZE = intsizeof(xm->feature.max_dom_size);
//	MAX_DOM_SIZE = xm->feature.max_dom_size;
//	VS_SIZE = xm->feature.vs_size;
//	CS_SIZE = xm->feature.cs_size;
//	BITDOMS_INTSIZE = BITDOM_INTSIZE * VS_SIZE;
//	BITSUP_INTSIZE = MAX_DOM_SIZE * BITDOM_INTSIZE;
//	BITSUPS_INTSIZE = BITSUP_INTSIZE * CS_SIZE;
//	BITSUBDOMS_INTSIZE = VS_SIZE * MAX_DOM_SIZE * BITDOMS_INTSIZE;
//	SUBCON_SIZE = VS_SIZE * MAX_DOM_SIZE * CS_SIZE;
//#pragma endregion 㳣
//#pragma region ԼϢ
//	//cudaMallocManaged(&vars_size, sizeof(int) * VS_SIZE);
//	//// ʼС
//	//for (int i = 0; i < xm->feature.vs_size; ++i)
//	//{
//	//	XVar* v = xm->vars[i];
//	//	XDom* d = xm->doms[v->dom_id];
//	//	vars_size[i] = d->size;
//	//}
//
//	// ʼscope
//	cudaMallocManaged(&scope, sizeof(int3) * CS_SIZE);
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		XCon *c = xm->cons[i];
//		scope[i].x = c->scope[0];
//		scope[i].y = c->scope[1];
//		scope[i].z = c->id;
//	}
//
//	////ʾ
//	//for (int i = 0; i < CS_SIZE; ++i)
//	//{
//	//	printf("scope[%d] = {%d, %d}\n", scope[i].z, scope[i].x, scope[i].y);
//	//}
//#pragma endregion ԼϢ
//#pragma region bitDom
//	cudaMallocManaged(&bitDom, sizeof(u32) * BITDOMS_INTSIZE);
//	cudaMallocManaged(&M_VarPre, sizeof(int) * VS_SIZE);
//
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		XVar* v = xm->vars[i];
//		XDom* d = xm->doms[v->dom_id];
//		const int dom_size = d->size;
//		// ǰʵINT
//		const int dom_int_size = intsizeof(dom_size);
//
//		for (int j = 0; j < BITDOM_INTSIZE; ++j)
//		{
//			const int idx = GetBitDomIndex(i, j);
//			//printf("idx = %d\n", idx);
//			// 
//			if (j < dom_int_size - 1)
//				bitDom[idx] = UINT32_MAX;
//			else if (j == dom_int_size - 1)
//				bitDom[idx] = UINT32_MAX << GetOffSet(dom_size);
//			else
//				bitDom[idx] = 0;
//		}
//
//		M_VarPre[i] = 1;
//	}
//
//	//for (int i = 0; i < VS_SIZE; ++i)
//	//{
//	//	for (int j = 0; j < BITDOM_INTSIZE; ++j)
//	//	{
//	//		int idx = GetBitDomIndex(i, j);
//	//		printf("var = %d, j = %d, idx = %d, bitDom = %x, pre= %x\n", i, j, idx, bitDom[idx], M_VarPre[i]);
//	//	}
//	//}
//#pragma endregion bitDom
//#pragma region bitSubDom
//	cudaMallocManaged(&bitSubDom, sizeof(u32)*BITDOMS_INTSIZE*VS_SIZE*MAX_DOM_SIZE);
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			const int start_idx = GetBitSubDomStartIndex(i, j);
//			for (int k = 0; k < BITSUBDOMS_INTSIZE; ++k)
//				bitSubDom[start_idx + k] = bitDom[k];
//			//(i,j)bitDom ĵ
//			//ȡi,j,iʼַ
//			const int ijistart = start_idx + i*BITDOM_INTSIZE;
//			for (int k = 0; k < BITDOM_INTSIZE; ++k)
//				// jKķΧ:j/32,j%32λΪ1
//				if (k == j >> U32_POS)
//					bitSubDom[ijistart + k] = U32_MASK1[j&U32_MOD_MASK];
//			//λΪ0
//				else
//					bitSubDom[ijistart + k] = 0;
//		}
//	}
//
//	//for (int i = 0; i < VS_SIZE; ++i)
//	//{
//	//	for (int j = 0; j < MAX_DOM_SIZE; ++j)
//	//	{
//	//		printf("sub problem:(%d, %d): ", i, j);
//	//		const int start_idx = GetBitSubDomStartIndex(i, j);
//	//		for (int k = 0; k < BITDOMS_INTSIZE; ++k)
//	//		{
//	//			printf("%x ", bitSubDom[start_idx + k]);
//	//		}
//	//		printf("\n");
//	//	}
//	//}
//#pragma endregion bitSubDom
//#pragma region bitSup
//	cudaMallocManaged(&bitSup, sizeof(uint2) * BITSUPS_INTSIZE);
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		XCon* c = xm->cons[i];
//		XRel* r = xm->rels[c->rel_id];
//		XVar* v[2] = { xm->vars[c->scope[0]], xm->vars[c->scope[1]] };
//		XDom* d[2] = { xm->doms[v[0]->dom_id], xm->doms[v[1]->dom_id] };
//
//		//ʼλ
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			for (int k = 0; k < BITDOM_INTSIZE; ++k)
//			{
//				const int idx = GetBitSupIndexByINTPrstn(c->id, j, k);
//				if (j < d[0]->size && (k < (d[1]->size >> U32_POS)))
//				{
//					//֧ȡ0x0000..., ͻȡ0xFFF...
//					bitSup[idx].x = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].y = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//				}
//				else if (k == (d[1]->size >> U32_POS))
//				{
//					bitSup[idx].x = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].y = (r->sem == SEM_CONFLICT) ? UINT32_MAX : 0;
//					bitSup[idx].x <<= U32_BIT - (d[1]->size & U32_MOD_MASK);
//					bitSup[idx].y <<= U32_BIT - (d[1]->size & U32_MOD_MASK);
//				}
//				else
//				{
//					bitSup[idx].x = 0;
//					bitSup[idx].y = 0;
//				}
//			}
//		}
//		//λֵ
//		for (int j = 0; j < r->size; ++j)
//		{
//			const int2 t = make_int2(r->tuples[j][0], r->tuples[j][1]);
//			//printf("c_id= %d, %d, %d\n", c->id, t.x, t.y);
//			const int2 idx = GetBitSupIndexByTuple(c->id, t);
//			//printf("idx = %d, %d\n", idx.x, idx.y);
//			if (r->sem == SEM_SUPPORT)
//			{
//				bitSup[idx.x].x |= U32_MASK1[t.y & U32_MOD_MASK];
//				bitSup[idx.y].y |= U32_MASK1[t.x & U32_MOD_MASK];
//			}
//			else
//			{
//				bitSup[idx.x].x &= U32_MASK0[t.y & U32_MOD_MASK];
//				bitSup[idx.y].y &= U32_MASK0[t.x & U32_MOD_MASK];
//			}
//		}
//		//// ʼλ
//		//for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		//{
//		//	printf("c_id = %d, j = %d: ", i, j);
//		//	for (int k = 0; k < BITDOM_INTSIZE; ++k)
//		//	{
//		//		const int idx = GetBitSupIndexByINTPrstn(c->id, j, k);
//		//		printf("%x, %x", bitSup[idx].x, bitSup[idx].y);
//		//	}
//		//	printf("\n");
//		//}
//	}
//#pragma endregion bitSup
//#pragma region Լ
//	cudaMallocManaged(&M_Con, sizeof(int)*CS_SIZE);
//	cudaMallocManaged(&M_ConEvt, sizeof(int) * CS_SIZE);
//	cudaMallocManaged(&M_ConPre, sizeof(int)*CS_SIZE);
//
//	for (int i = 0; i < CS_SIZE; ++i)
//	{
//		M_Con[i] = i;
//		M_ConEvt[i] = i;
//		M_ConPre[i] = 1;
//		//printf("i = %d , M_Con = %d, M_ConEvt = %d, M_ConPre = %d\n", i, M_Con[i], M_ConEvt[i], M_ConPre[i]);
//	}
//#pragma endregion Լ
//#pragma region Լ
//	cudaMallocManaged(&S_ConPre, sizeof(int)*SUBCON_SIZE);
//	cudaMallocManaged(&S_ConEvt, sizeof(int3)*SUBCON_SIZE);
//	cudaMallocManaged(&S_Con, sizeof(int3)*SUBCON_SIZE);
//	cudaMallocManaged(&S_VarPre, sizeof(int)*VS_SIZE*MAX_DOM_SIZE*VS_SIZE);
//	cudaMallocManaged(&S_Var, sizeof(int3)*VS_SIZE*MAX_DOM_SIZE*VS_SIZE);
//
//	for (int i = 0; i < VS_SIZE; ++i)
//	{
//		const XVar* v = xm->vars[i];
//		for (int j = 0; j < MAX_DOM_SIZE; ++j)
//		{
//			for (int k = 0; k < CS_SIZE; ++k)
//			{
//				// (i, j) kΪԼid
//				const int idx = (i*MAX_DOM_SIZE + j)*CS_SIZE + k;
//				//i*xm->feature.max_dom_size*xm->feature.cs_size + j*xm->feature.cs_size + k;
//
//				S_Con[idx].x = i;
//				S_Con[idx].y = j;
//				S_Con[idx].z = k;
//
//				S_ConEvt[idx].x = i;
//				S_ConEvt[idx].y = j;
//				S_ConEvt[idx].z = k;
//
//				S_ConPre[idx] = 1;
//				//printf("S_Con = (%d, %d, %d), S_ConEvt = (%d, %d, %d), pre = %d\n", S_Con[idx].x, S_Con[idx].y, S_Con[idx].z, S_ConEvt[idx].x, S_ConEvt[idx].y, S_ConEvt[idx].z, S_ConPre[idx]);
//			}
//
//			for (int k = 0; k < VS_SIZE; ++k)
//			{
//				//(i, j) kΪid
//				const int idx = (i*MAX_DOM_SIZE + j)*VS_SIZE + k;
//				S_Var[idx].x = i;
//				S_Var[idx].y = j;
//				S_Var[idx].z = k;
//
//				S_VarPre[idx] = 1;
//
//				//printf("S_Var = (%d, %d, %d), S_VarPre = %d\n", S_Var[idx].x, S_Var[idx].y, S_Var[idx].z, S_VarPre[idx]);
//			}
//		}
//	}
//#pragma endregion Լ
//
//#pragma region й
//	//ѹBLOCK
//	MCC_BLOCK = GetTopNum(CS_SIZE, num_threads);
//	MCC_BCount.resize(MCC_BLOCK, 0);
//	MCC_BOffset.resize(MCC_BLOCK, 0);
//	MCC_BlocksCount = thrust::raw_pointer_cast(MCC_BCount.data());
//	MCC_BlocksOffset = thrust::raw_pointer_cast(MCC_BOffset.data());
//#pragma endregion
//
//}
//
//__global__ void ConCheckMain(int* ConEvt, int* btSp, int2* scp)
//{
//	const int c_id = blockIdx.x;
//	//ȡԼbitSupĿʼ
//	const int start_idx = GetBitSupIndexById(c_id);
//	const int2 sp = scp[c_id];
//	extern __shared__ int2[];
//}
//
//void ConstraintCheckMain()
//{
//	//num_threadsÿɱ
//	ConCheckMain << <CS_SIZE, num_threads >> >();
//}
//
//float SACGPU()
//{
//	//1. ִAC
//	//1.1. ѹ
//	CompactQueueMain();
//	//1.2. Լ
//	ConstraintCheckMain();
//}
//
//void DelGPUModel()
//{
//	cudaFree(scope);
//	cudaFree(bitDom);
//	cudaFree(M_VarPre);
//	cudaFree(bitSubDom);
//	cudaFree(bitSup);
//	cudaFree(M_Con);
//	cudaFree(M_ConEvt);
//	cudaFree(M_ConPre);
//	cudaFree(S_ConPre);
//	cudaFree(S_ConEvt);
//	cudaFree(S_Con);
//	cudaFree(S_Var);
//	cudaFree(S_VarPre);
//}
//
//

	.text
	.file	"cudasac.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
