I don't know this error.

Hello, I’m an OpenCL beginner.

I’d like to develop of hextile algorithm.
I studied sample source, article, and another comment.
But I don’t know why are they occur error!!
I’m very stressed out, because this problem occured last week.

Please, help me to solve this problem.

Error message is…
[ul]ERROR : Failed to build program executable!
ptxas ptx input, line 59; fatal : Parsing error near ‘;’: syntax error
: Retrieving binary for ‘anonymous_jit_identity’, for gpu=‘sm_12’, usage mode=’[/ul]

There seems to be some kind of syntax error in your kernel code. We could have a look at it if you post it.

source code is that…

__kernel __attribute__((vec_type_hint(int)))
void mysubrectEncode(
					__global int* destoffset,	// out data
					__global ulong *dest,	// out data
					__global ulong *subrectData,// in data
					const int w, 
					const int h, 
					ulong bg,  
					int mono) 	
{
	int yH = get_group_id(0);	
	int xW = get_local_id(1);
	int addr = yH*w;

	ulong cl;
	int i,j;
	int hx = 0, vx = 0, hy, vy;
	int hyflag;
	int hw,hh,vw,vh;
	int thex, they, thew, theh;
	int numsubs = 0;
	int newLen;
	int rectoffset;

	__global ulong4 *line;
	__global ulong4 *seg;
destoffset = 0;
	rectoffset = *destoffset;
	*destoffset += 1;

	line = (subrectData + (yH*w));

			if(bg != (ulong)(line+xW))
			{
				cl = (ulong)line+xW;
				hy = yH-1;
				hyflag = 1;

				for(j=yH; j<h; j++)
				{
BREAK:					
					seg = (__global ulong4 *)(subrectData + (j*w));	

					if(cl != (ulong)(seg + xW))
					{ 
						goto BREAK;
						//break;
					}
					i = xW;

					while((cl == (ulong)(seg+i)) && (i<w))
						i += 1;
						
					i -= 1;


					if(j==yH)
						vx = hx = i;
					if(i<vx)
						vx = i;

					if((hyflag > 0) && (i >= hx))
					{
						hy += 1;
					}
					else
					{
						hyflag = 0;
					}
				}
........
}

I don’t know…why syntax error…

I could not find any misplaced ‘,’ near line 59, which was what the error message hinted at. I have never tried using gotos in kernel code, try replacing it with a ‘continue;’

I could not find any misplace, too.
This line is changed, and then execute build.
After a moment, the error message is continue…
(error line is not changed any-number… :?: )

Someone said, you aren’t control the memory pointer… or pointer conversions problem.

But I don’t know why occur this error message.

I used global memory space, the other space is not used…Because at the part of memory is division each others.(I think that…)

I’m very sad… -_?

What is your OpenCL platform and how do you compiile this kernel?

I compile this kernel code is that…

shrLog(LOGBOTH, 0.0, "Initializing OpenCL
");
		oclHandles.context = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, &err);
		shrCheckError(err, CL_SUCCESS);

		err = clGetContextInfo(oclHandles.context, CL_CONTEXT_DEVICES, 0, NULL, &dataByte);
		oclHandles.devices = (cl_device_id*)malloc(dataByte);
		
		err |= clGetContextInfo(oclHandles.context, CL_CONTEXT_DEVICES, dataByte, oclHandles.devices, NULL);
		shrCheckError(err, CL_SUCCESS);

		GetAndLogDevice(oclHandles.devices, argc, &argv, &oclHandles.context);	// print device data

		//Create a command-queue
		oclHandles.queue = clCreateCommandQueue(oclHandles.context, oclHandles.devices[0], 0, &err);
		shrCheckError(err, CL_SUCCESS);

		// device Info
		size_t returned_size = 0;
		size_t max_workgroup_size = 0;
		err = clGetDeviceInfo(*oclHandles.devices, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &max_workgroup_size, &returned_size);
		if (err != CL_SUCCESS)
		{
			printf("Error: Failed to retrieve device info!
");
			return EXIT_FAILURE;
		}
// 3. allocating inputCL buffer..
	shrLog(LOGBOTH, 0.0, "...allocating inputCL buffers
");
		

		// if you use CL_MEM_COPY_HOST_PTR attibute, you must write 4 parameter(void*) data.
		// Another attribute - you choose that 4 parameter write or NULL.
		clBuffers.inputCL = clCreateBuffer(oclHandles.context, CL_MEM_READ_WRITE, compressedSize, NULL, &err);
		if(!clBuffers.inputCL)
		{
			printf("ERROR : Failed to allocating inputCL buffer!
");
			return EXIT_FAILURE;
		}
		// Hextile ??? ?? ?? raw ??? ? ??? ??? size? image? ???? ??
		clBuffers.outputCL = clCreateBuffer(oclHandles.context, CL_MEM_WRITE_ONLY, compressedSize, NULL, &err);
		if(!clBuffers.outputCL)
		{
			printf("ERROR : Failed to allocating OutputCL buffer!
");
			return EXIT_FAILURE;
		}
//4. Program Setup
	//
	shrLog(LOGBOTH, 0.0, "Initializing OpenCL Hextile...
");
		// 		ret = InitOpenCL(clGPUContext, clCommandQueue, argv);
		size_t program_length;
		cl_device_id device_id;

		// kernel.cl loading
		shrLog(LOGBOTH, 0.0, "..loading Kernel.cl
");
		char* source_path = shrFindFilePath("SubRectEncode.cl", argv[0]);
		shrCheckError(source_path != NULL, shrTRUE);
//		const char* filename = "SubRectEncode.cl";
//		char *source_path = load_program_source(filename);

		char *kernel_source = oclLoadProgSource(source_path, "", &program_length);
		//shrCheckError(kernel_source != NULL, shrTRUE);

		// Connect to a GPU computing device
		err = clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
		if(err != CL_SUCCESS)
		{
			printf("ERROR : Failed to create a device group!
");
			return EXIT_FAILURE;
		}

		// create kernel program
		shrLog(LOGBOTH, 0.0, "..creating kernel program
");
		oclHandles.program = clCreateProgramWithSource(oclHandles.context, 1, (const char **)&kernel_source, &program_length, &err);
		if(!oclHandles.program || err != CL_SUCCESS)
		{
			printf("ERROR : Failed to create computing program!!
");
			return EXIT_FAILURE;
		}

		
		// build kernel program
		shrLog(LOGBOTH, 0.0, "...building kernel program
");

		err = clBuildProgram(oclHandles.program, 0, NULL, NULL, NULL, NULL); // this erro!
		if(err != CL_SUCCESS)
		{
			size_t len;
			char buffer[4096];

			printf("ERROR : Failed to build program executable!
");
			clGetProgramBuildInfo(oclHandles.program, device_id, CL_PROGRAM_BUILD_LOG, 
				sizeof(buffer), buffer, &len);
			printf("%s
", buffer);
			return EXIT_FAILURE;
		}

I write the error statement up to the code.

Oh, I have one question…
The sample source code at NVIDIA are not include the clGetDeviceIDs() function.
But another sample at web-stie is included the clGetDeviceIDs() function.
Is it important to me??

please, advise to me.

I don’t know why is error.

ptxas ptx input, line 59; fatal : Parsing error near ‘;’: syntax error

is the error. You could try to comment out the line 59 (and related) to see, if this “solve” the error.
When I count correctly, there is the problem :


               if(j==yH)
                  vx = hx = i;

If so, try to use


               if(j==yH)
               {
                  vx = i;
                  hx = i;
                }

first, thank you very much your advise.

I fixed it today.

But the error is still existence…

I need to another advise. Pleaz tell me more. :slight_smile:

It does not seem to be a quick and easy fix anyone can give you. I would suggest that you make a copy of some example you have that is working and incrementally turn it into what you really want. Eventually you will either add the piece of code that causes the error and then perhaps easier find a fix, or you will unknowingly do something better in the second version and not have the error at all.

The problem is not directly related to your code, but rather to NVIDIA’s OpenCL implementation which is built upon CUDA. I have seen similar error messages.

NVIDIA uses a modified version of the clang compiler and LLVM to bring the OpenCL C code into a intermediate format. While your code is syntactically correct, the output of the clang compiler is not. There is nothing that you can do, except waiting for an OpenCL implementation that is not broken. Try to use a more recent version if you can.

Alternative solutions are:

  • use CUDA
  • use ATI/AMD hardware and their OpenCL implementation
  • redesign your code and pray

I using your suggestion that I make copy a piece of my code. And then I get my error but also I don’t understand with this error.

This error is occur in ‘for roop’.

 __kernel __attribute__((vec_type_hint(int)))
void subrectEncode32(
					__global int* resDestoffset,	// out data
					__global ulong *dest,		// out data
					__global ulong *subrectData,// in data
					const int w, 
					const int h, 
					ulong bg,  
					int mono) 	//bool mono)
{
	int result = get_global_id(0);	// width * height? ??? ??
	
	int yH = get_group_id(0);	// height? ?? ???
	int xW = get_local_id(0);	// width? ?? ??? 
	
	ulong cl;
	int i,j;
	int hx = 0;
	int vx = 0;
	int hy;
	int vy;
	int hyflag;
	int hw, hh, vw, vh;
	int thex, they, thew, theh;
	int numsubs = 0;
	int newLen;
	int rectoffset;
	
	int addr = yH*w;
	
	__global ulong *line;
	__global ulong *seg;
	
//	ulong4 thing;
	int destoffset;

	destoffset = 0;
	rectoffset = destoffset;
	destoffset = destoffset + 1;

	line = (subrectData + (yH*w));

	if(bg != (ulong)(line+xW))
	{
		cl = (ulong)line+xW;
		hy = yH-1;
		hyflag = 1;

		for(j=yH; j<h; j++)    // error is here! 
		{
		     .......


the error is also the same before.
“ptxas ptx input, line 99; error : Arguments mismatch for instruction ‘mov’
: Retrieving binary for ‘anonymous_jit_identity’, for gpu=‘sm_12’, usage mode=’”