Skip to content

Why program.build() always says "}"didn't match? #41

New issue
New issue
Open
Open
Why program.build() always says "}"didn't match?#41
@liushaohuai5

Description

@liushaohuai5
liushaohuai5
opened on Jun 24, 2019

:441:12: error: expected '}'
:439:10: note: to match this '{'
{
^

kernel file:

// TODO: Add OpenCL kernel code here.

__kernel void dpqc(
int Ndim,
//__global float *G,
__constant float *G,
//__global float *B,
__constant float *B,
__global float *e,
__global float *f,
__global float *dP,
__global float *dQ,
__global float *temp1,
__global float *temp2,
__constant float *Pc,
__constant float *Qc,
__constant float *Uc,
__global float *dU2,
int Mdim,
__global float *P,
__global float *Q,
__global int *PQV_flag){

			int k;
			int i = get_global_id(0);
			float tmp1, tmp2;
			printf("%d\n",PQV_flag[i]);
			if(PQV_flag[i]==1){		// 1~Mdim是只算PQ节点
				tmp1 = 0.0;
				tmp2 = 0.0;
				#pragma unroll
				for(k=0;k<Ndim;k++){ 
					tmp1 += G[i*Ndim+k]*e[k] - B[i*Ndim+k]*f[k];
					tmp2 += G[i*Ndim+k]*f[k] + B[i*Ndim+k]*e[k];
					//printf("%f\n",tmp2);
				}
				//printf("%f\n",tmp2);
				P[i] = (e[i]*tmp1 + f[i]*tmp2);
				dP[i] = Pc[i] - (e[i]*tmp1 + f[i]*tmp2);
				Q[i] = f[i]*tmp1 - e[i]*tmp2;
				dQ[i] = Qc[i] - f[i]*tmp1 + e[i]*tmp2;
				//dU2[i] = Uc[i]*Uc[i] - (e[i]*e[i] + f[i]*f[i]); 
				temp1[i] = tmp1;		//中间量，其他内核可以用到
				temp2[i] = tmp2;
			}
			else if(PQV_flag[i]==0){	// Mdim+1~Ndim-1是只算PV节点
				tmp1 = 0.0;
				tmp2 = 0.0;
				#pragma unroll
				for(k=0;k<Ndim;k++){ 
					tmp1 += G[i*Ndim+k]*e[k] - B[i*Ndim+k]*f[k];
					tmp2 += G[i*Ndim+k]*f[k] + B[i*Ndim+k]*e[k];
					//printf("%f\n",tmp1);
				}
				P[i] = (e[i]*tmp1 + f[i]*tmp2);
				dP[i] = Pc[i] - (e[i]*tmp1 + f[i]*tmp2);
				//printf("%f %f %f\n",Uc[i], e[i], f[i]);
				dU2[i] = Uc[i]*Uc[i] - (e[i]*e[i] + f[i]*f[i]); 
				//printf("%d %f\n", i ,dU2[i]);
				temp1[i] = tmp1;		//中间量，其他内核可以用到
				temp2[i] = tmp2;
			}
		}

__kernel void jmcc(
int M,
int N,
//__global float *G,
__constant float *G,
//__global float *B,
__constant float *B,
__global float *e,
__global float *f,
__global float *temp1,
__global float *temp2,
__global float *J,
__global int *PQV_flag){

				int k;
				int Jdim;
				Jdim = (2*N-2)*2;
				float tmp1,tmp2;
				int i = get_global_id(0);
				if(PQV_flag[i]==1){ 
					for(k=0; k<N-1; k++){ 
						if(i != k){ //单次计算H,N,J,L，Jaccobi矩阵一行2*N-2个元素
							J[i*Jdim+k*2] = -(G[i*N+k]*e[i] + B[i*N+k]*f[i]);	//Hij
							J[i*Jdim+k*2+1] = B[i*N+k]*e[i] - G[i*N+k]*f[i];	//Nij = Bij*ei - Gij*fi
							J[i*Jdim+2*k+(2*N-2)] = J[i*Jdim+k*2+1];			//Jij = Nij
							J[i*Jdim+2*k+1+(2*N-2)] = -J[i*Jdim+k*2];			//Lij = -Hij
							//printf("%f %f %f %f")	
						}
						else{ 
							tmp1=0;
							tmp2=0;
							for(int j=0;j<N;j++){ 
								tmp1 += G[i*N+j]*e[j] - B[i*N+j]*f[j];
								tmp2 += G[i*N+j]*f[j] + B[i*N+j]*e[j];
								//printf("%f\n",tmp1);
							}
							J[i*Jdim+k*2] =  -tmp1 - G[i*N+k]*e[i] - B[i*N+k]*f[i];				//Hii
							J[i*Jdim+k*2+1] = -tmp2 + B[i*N+k]*e[i] - G[i*N+k]*f[i];			//Nii
							J[i*Jdim+2*k+(2*N-2)] = tmp2 + B[i*N+k]*e[i] - G[i*N+k]*f[i];		//Jii
							J[i*Jdim+2*k+1+(2*N-2)] = -tmp1 + G[i*N+k]*e[i] + B[i*N+k]*f[i];	//Lii
							
						}
					}
				}
				else if(PQV_flag[i]==0){ 
					for(k=0; k<N-1; k++){ 
						if(i != k){
							J[i*Jdim+k*2] = -(G[i*N+k]*e[i] + B[i*N+k]*f[i]);	//Hij
							J[i*Jdim+k*2+1] = B[i*N+k]*e[i] - G[i*N+k]*f[i];	//Nij = Bij*ei - Gij*fi
							J[i*Jdim+2*k+(2*N-2)] = 0;							//Rij
							J[i*Jdim+2*k+1+(2*N-2)] = 0;						//Sij
							
						}
						else{ 
							tmp1=0;
							tmp2=0;
							for(int j=0;j<N;j++){ 
								tmp1 += G[i*N+j]*e[j] - B[i*N+j]*f[j];
								tmp2 += G[i*N+j]*f[j] + B[i*N+j]*e[j];
								//printf("%f\n",tmp1);
							}
							J[i*Jdim+k*2] =  -tmp1 - G[i*N+k]*e[i] - B[i*N+k]*f[i];			//Hii
							J[i*Jdim+k*2+1] = -tmp2 + B[i*N+k]*e[i] - G[i*N+k]*f[i];			//Nii
							J[i*Jdim+2*k+(2*N-2)] = -2*e[i];
							J[i*Jdim+2*k+1+(2*N-2)] = -2*f[i];
							
						}
					}
				}	
			}

__kernel void transpose(
__global float input_mat,
__global float output_mat,
int width,
int height)
{
int row = get_global_id(0);
//int col = get_global_id(1);
for(int col=0; col<width; col++){
output_mat[rowwidth+col] = input_mat[colheight+row];
}
}

__kernel void qr(
__local float *u_vec,
__global float *a_mat,
__global float *q_mat,
__global float *p_mat,
__global float *prod_mat) {

				local float u_length_squared, dot;
				float prod, vec_length = 0.0f;

				int id = get_local_id(0);
				int num_cols = get_global_size(0);

				/* Load first column into local memory as u vector */
				u_vec[id] = a_mat[id*num_cols];
				//printf("%d %f\n",id,u_vec[id]);
				barrier(CLK_LOCAL_MEM_FENCE);

				/* Find length of first A column and u vector */
				if(id == 0) {
					for(int i=1; i<num_cols; i++) {
						vec_length += u_vec[i] * u_vec[i];
					}
					u_length_squared = vec_length;
  
					vec_length = sqrt(vec_length + u_vec[0] * u_vec[0]);
					a_mat[0] = vec_length;
					u_vec[0] -= vec_length;
					u_length_squared += u_vec[0] * u_vec[0];
				}
				else {
					a_mat[id*num_cols] = 0.0f;
				}
				barrier(CLK_GLOBAL_MEM_FENCE);

				/* Transform further columns of A */
				for(int i=1; i<num_cols; i++) {
					dot = 0.0f;
					if(id == 0) {
						for(int j=0; j<num_cols; j++) {
							 dot += a_mat[j*num_cols + i] * u_vec[j];
						}
					}
					barrier(CLK_LOCAL_MEM_FENCE);
					a_mat[id*num_cols + i] -= 2 * u_vec[id] * dot / u_length_squared;
				}

				/* Update Q matrix */
			   for(int i=0; i<num_cols; i++) {
				  q_mat[id*num_cols + i] = -2 * u_vec[i] * 
						u_vec[id] / u_length_squared;
			   }
			   q_mat[id*num_cols + id] += 1;
			   barrier(CLK_GLOBAL_MEM_FENCE); 

			   /* Loop through other columns */
			   for(int col = 1; col < num_cols-1; col++) {

				  /* Load new column into memory */
				  u_vec[id] = a_mat[id * num_cols + col];
				  barrier(CLK_LOCAL_MEM_FENCE);

				  /* Find length of A column and u vector */
				  if(id == col) {
						vec_length = 0.0f;
						for(int i = col + 1; i < num_cols; i++) {
							vec_length += u_vec[i] * u_vec[i];
						}
						u_length_squared = vec_length;
						vec_length = sqrt(vec_length + u_vec[col] * u_vec[col]);
						u_vec[col] -= vec_length;
						u_length_squared += u_vec[col] * u_vec[col];
						a_mat[col * num_cols + col] = vec_length;
				  }
				  else if(id > col) {
					a_mat[id * num_cols + col] = 0.0f;
					}
					barrier(CLK_GLOBAL_MEM_FENCE);

				/* Transform further columns of A */
				  for(int i = col+1; i < num_cols; i++) {
					 if(id == 0) {
						dot = 0.0f;
						for(int j=col; j<num_cols; j++) {
						   dot += a_mat[j*num_cols + i] * u_vec[j];
						}
					 }
					 barrier(CLK_LOCAL_MEM_FENCE);
     
					 if(id >= col)
						a_mat[id*num_cols + i] -= 2 * u_vec[id] * 
							  dot / u_length_squared;
					 barrier(CLK_GLOBAL_MEM_FENCE);
				  }

				  /* Update P matrix */
				  if(id >= col) {
					 for(int i=col; i<num_cols; i++) {
						p_mat[id*num_cols + i] = -2 * u_vec[i] * 
							  u_vec[id] / u_length_squared;
					 }
					 p_mat[id*num_cols + id] += 1;
				  }
				  barrier(CLK_GLOBAL_MEM_FENCE); 

				  /* Multiply q_mat * p_mat = prod_mat */
				  for(int i=col; i<num_cols; i++) {
					 prod = 0.0f;
					 for(int j=col; j<num_cols; j++) {
						prod += q_mat[id*num_cols + j] * p_mat[j*num_cols + i];
					 }     
					 prod_mat[id*num_cols + i] = prod;  
				  }
				  barrier(CLK_GLOBAL_MEM_FENCE); 

				  /* Place the content of prod_mat in q_mat */
				  for(int i=col; i<num_cols; i++) {
					 q_mat[id*num_cols + i] = prod_mat[id*num_cols + i];
				  }
				  barrier(CLK_GLOBAL_MEM_FENCE); 
			   }
			}

__kernel void sevc1(
int dim,
__global float *R,
__global float Q,
__global float b,
__global float x,
__local float tmp
){
float tmp1=0;
int id = get_global_id(0);
for(int i=0; i<dim; i++){
tmp[id] += Q[iddim+i] * b[i];
}
barrier(CLK_GLOBAL_MEM_FENCE);
//printf("b'[%d]=%f\n",id,tmp[id]);
for(int i=dim-1; i>=0; i--){
if(i==dim-1){
x[i] = tmp[i]/R[idim+i]; //x(5) = b'(5)/R[5][5];
//printf("x[%d] = %f %f\n",i,x[i],R[idim+i]);
}
else{
for(int j=i+1; j<dim; j++){
//if(i<dim-1){
tmp1 += R[idim+j]x[j];
//}
//else{
//tmp1=0;
//}
}
x[i] = (tmp[i]-tmp1)/R[idim+i];
tmp1=0;
}
}
}

			__kernel void dpqc_sparse(
						int Ndim,
						__constant float *G_data,
						__constant int *G_ICOL,
						__constant int *G_ICFR,
						__constant float *B_data,
						__constant int *B_ICOL,
						__constant int *B_ICFR,
						__global float *e,
						__global float *f,
						__global float *dP,
						__global float *dQ,
						__constant float *Pc,
						__constant float *Qc,
						__constant float *Uc,
						__global float *dU2,
						int Mdim,
						__global float *P,
						__global float *Q,
						__global int *PQV_flag){
						
						int i = get_global_id(0);
						//int k;
						float tmp1=0.0, tmp2=0.0, tmp3=0.0, tmp4=0.0;
						//printf("%d %f\n",i,G_data[i]);
						//printf("%d\n",PQV_flag[i]);
						if(/*i<Mdim*/PQV_flag[i]==1){
							tmp1=0;tmp2=0;tmp3=0;tmp4=0;
							for(int j = G_ICFR[i]; j < G_ICFR[i+1]; j++){
								tmp1 += G_data[j] * e[G_ICOL[j]];		//G*e
								tmp2 += G_data[j] * f[G_ICOL[j]];		//G*f
							}
							for(int j = B_ICFR[i]; j < B_ICFR[i+1]; j++){
								tmp3 += B_data[j] * e[B_ICOL[j]];		//B*e
								tmp4 += B_data[j] * f[B_ICOL[j]];		//B*f
							}
							dP[i] = Pc[i] - (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
							dQ[i] = Qc[i] - (f[i]*(tmp1-tmp4) - e[i]*(tmp2+tmp3));
							P[i] = (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
							Q[i] = (f[i]*(tmp1-tmp4) - e[i]*(tmp2+tmp3));
						}
						else if(/*i<Ndim-1*/PQV_flag[i]==0){
							tmp1=0;tmp2=0;tmp3=0;tmp4=0;
							for(int j = G_ICFR[i]; j < G_ICFR[i+1]; j++){
								tmp1 += G_data[j] * e[G_ICOL[j]];		//G*e
								tmp2 += G_data[j] * f[G_ICOL[j]];		//G*f
							}
							for(int j = B_ICFR[i]; j < B_ICFR[i+1]; j++){
								tmp3 += B_data[j] * e[B_ICOL[j]];		//B*e
								tmp4 += B_data[j] * f[B_ICOL[j]];		//B*f
							}

							dP[i] = Pc[i] - (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
							//dQ[i] = Qc[i] - (f[i]*(tmp1-tmp4) - e[i]*(tmp2+tmp3));
							//printf("%f %f %f\n",Uc[i], e[i], f[i]);
							dU2[i] = Uc[i]*Uc[i] - (e[i]*e[i] + f[i]*f[i]);
							P[i] = (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
						}
						else if(PQV_flag[i] == 2){
							tmp1=0;tmp2=0;tmp3=0;tmp4=0;
							for(int j = G_ICFR[i]; j < G_ICFR[i+1]; j++){
								tmp1 += G_data[j] * e[G_ICOL[j]];		//G*e
								tmp2 += G_data[j] * f[G_ICOL[j]];		//G*f
							}
							for(int j = B_ICFR[i]; j < B_ICFR[i+1]; j++){
								tmp3 += B_data[j] * e[B_ICOL[j]];		//B*e
								tmp4 += B_data[j] * f[B_ICOL[j]];		//B*f
							}
							dP[i] = Pc[i] - (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
							dQ[i] = Qc[i] - (f[i]*(tmp1-tmp4) - e[i]*(tmp2+tmp3));
							P[i] = (e[i]*(tmp1-tmp4) + f[i]*(tmp2+tmp3));
							Q[i] = (f[i]*(tmp1-tmp4) - e[i]*(tmp2+tmp3));
						}
					}

__kernel void jmcc_sparse(
int M,
int N,
__constant float *G_data,
__constant float *B_data,
__constant float *ICFR,
__constant int *ROW,
__constant int *COL,
int num,
__global int *PQV_flag,
__global float *e,
__global float *f,
__global float *J){

				int Jdim = (2*N-2)*2;
				float tmp1,tmp2,tmp3,tmp4;
				int i = get_global_id(0);
				printf("%2d %d %d \n",i,G_data[ROW[i]],G_data[COL[i]]);		
					if( PQV_flag[i] == 1 ){		//PQV_flag[i] == 1,表示PQ节点
						for(int j=ICFR[i];j<ICFR[i+1];j++){		//排零运算
							if(i!=COL[j] && COL[j]<7){				//非对角块时
								J[ i*Jdim + 2*COL[j] ] = -(G_data[j]*e[i] + B_data[j]*f[i]);				//Hij
								J[ i*Jdim + 2*COL[j] + 1 ] = B_data[j]*e[i] - G_data[j]*f[i];					//Nij	
								J[ i*Jdim + (2*N-2) + 2*COL[j] ] = B_data[j]*e[i] - G_data[j]*f[i];		//Jij
								J[ i*Jdim + (2*N-2) + 2*COL[j] + 1] = G_data[j]*e[i] + B_data[j]*f[i];	//Lij
							}
							else if(i == COL[j] && COL[j]<7){
								tmp1=0;tmp2=0;
								for(int k=ICFR[i];k<ICFR[i+1];k++){
									tmp1 += G_data[k]*e[COL[k]] - B_data[k]*f[COL[k]];			//求和Gij*ej - Bij*fj
									tmp2 += G_data[k]*f[COL[k]] + B_data[k]*e[COL[k]];
								}
								J[ i*Jdim + 2*COL[j] ] = -tmp1 - G_data[j]*e[i] - B_data[j]*f[i];
								J[ i*Jdim + 2*COL[j] + 1 ] = -tmp2 - G_data[j]*f[i] + B_data[j]*e[i];
								J[ i*Jdim + (2*N-2) + 2*COL[j] ] = tmp2 - G_data[j]*f[i] + B_data[j]*e[i];
								J[ i*Jdim + (2*N-2) + 2*COL[j] + 1 ] = -tmp1 + G_data[j]*e[i] - B_data[j]*f[i];
							}
						}
					}
					else if( PQV_flag[i] == 0 ){	//PQV_flag[i] == 0,表示第i个节点从PQ节点变为了PV节点
						for(int j=ICFR[i];j<ICFR[i+1];j++){
							if(i!=COL[j] && COL[j]<7){
								J[ i*Jdim + 2*COL[j] ] = -(G_data[j]*e[i] + B_data[j]*f[i]);				//Hij
								J[ i*Jdim + 2*COL[j] + 1 ] = B_data[j]*e[i] - G_data[j]*f[i];			//Nij

							}
							else if(i == COL[j] && COL[j]<7){
								tmp1=0;tmp2=0;
								for(int k=ICFR[i];k<ICFR[i+1];k++){
									tmp1 += G_data[k]*e[COL[k]] - B_data[k]*f[COL[k]];			//求和Gij*ej - Bij*fj
									tmp2 += G_data[k]*f[COL[k]] + B_data[k]*e[COL[k]];
								}
								J[i*Jdim + 2*COL[j]] = -tmp1 - G_data[j]*e[i] - B_data[j]*f[i];
								J[i*Jdim + 2*COL[j]+1] = -tmp2 - G_data[j]*f[i] + B_data[j]*e[i];
								for(int m=0;m<(N-1);m++)
								{
									J[i*Jdim+(2*N-2)+2*m]=-2*e[i];
									J[i*Jdim+(2*N-2)+2*m+1]=-2*f[i];
								}
							}
						}
					}
				}

JAVA FILE:
package Flow_OpenCL;

import java.io.File;
import java.io.IOException;

import com.nativelibs4java.*;
import com.nativelibs4java.opencl.CLBuildException;
import com.nativelibs4java.opencl.CLContext;
import com.nativelibs4java.opencl.CLDevice;
import com.nativelibs4java.opencl.CLException;
import com.nativelibs4java.opencl.CLKernel;
import com.nativelibs4java.opencl.CLPlatform;
import com.nativelibs4java.opencl.CLProgram;
import com.nativelibs4java.opencl.CLQueue;
import com.nativelibs4java.opencl.JavaCL;
import com.nativelibs4java.util.IOUtils;

public class Flow1 {

public static void main(String[] args) throws Exception {
	
	CLPlatform[] Platforms = JavaCL.listPlatforms();					//查询可用平台
	CLDevice device = Platforms[0].getBestDevice();						//调用设备
	String platformName = Platforms[0].getName();						//获取平台名称
	String deviceName = device.getName();								//获取设备名称
	System.out.println("平台:"+platformName+" 设备:"+deviceName);		//打印平台和设备名称

	//CLContext context = JavaCL.createBestContext();		//为设备创建上下文
	CLContext context = JavaCL.createContext(null, device);
	//CLQueue queue = context.createDefaultQueue((CLDevice.QueueProperties[] )null);
	CLQueue queue = context.createDefaultProfilingQueue();
	
	
	//jmcc.cl-done
	//dpqc_sparse.cl-done
	//qr.cl-done
	//transpose.cl-done
	//sevc1.cl-done
	String kernelText = IOUtils.readText(new File("E:\\LSH\\JavaCL\\src\\Flow_OpenCL\\FLOW.cl"));
	CLProgram program = context.createProgram(kernelText);
	/*
	program.addInclude("E:\\LSH\\JavaCL\\src\\Flow_OpenCL\\qr.cl");
	String qrText = IOUtils.readText(new File("E:\\LSH\\JavaCL\\src\\Flow_OpenCL\\qr.cl"));
	program.addSource(qrText);
	*/
	program.build();
	
	CLKernel[] kernels = program.createKernels();
	System.out.println("完成");
	
	
}

}

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions