calling cl_arm_import_memory failed with error code -6 on RK3399

Kevin - thanks for your prompt response.

unfortunately, we are working a near-real-time video application runs on i.MX8 and the current CPU usage ratio is almost hit the cap. Hence we are trying to offload some computation (e.g., feature points tracking/matrix computation/HOG feature calculation) to the GPU. Initially we tried to use clCreateBuffer, but it looks like the gain of offloading the compute to the GPU is offset by the memory copy (data will need to be passed into GPU and result will need to be passed back to the CPU on frame basis) and consequently we cannot meet the real time requirement (33ms per frame). This is why we are trying to check whether zero-copy can help.

I reattached the code below, can you take another look? btw - what is error code -6?

#define WIDTH  100
#define HEIGHT 100

cl_platform_id platform;

status = clGetPlatformIDs(1, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, 
               CL_DEVICE_TYPE_GPU,
  		       1,
		       &device,
		       NULL);
context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);

char *allocptr = malloc(WIDTH*HEIGHT*2);

cl_mem buffer = clImportMemoryARM(context,
								  CL_MEM_READ_WRITE, 
								  NULL,
								  allocptr,
                                  WIDTH * HEIGHT * 2,
                                  &error );

if (error ==CL_SUCCESS)
{
	printf("sucess\n");
}
else
{
	printf("error %d\n.",error);
}

Kevin - thanks for your prompt response.

unfortunately, we are working a near-real-time video application runs on i.MX8 and the current CPU usage ratio is almost hit the cap. Hence we are trying to offload some computation (e.g., feature points tracking/matrix computation/HOG feature calculation) to the GPU. Initially we tried to use clCreateBuffer, but it looks like the gain of offloading the compute to the GPU is offset by the memory copy (data will need to be passed into GPU and result will need to be passed back to the CPU on frame basis) and consequently we cannot meet the real time requirement (33ms per frame). This is why we are trying to check whether zero-copy can help.

I reattached the code below, can you take another look? btw - what is error code -6?

#define WIDTH  100
#define HEIGHT 100

cl_platform_id platform;

status = clGetPlatformIDs(1, &platform, NULL);
cl_device_id device;
clGetDeviceIDs(platform, 
               CL_DEVICE_TYPE_GPU,
  		       1,
		       &device,
		       NULL);
context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, NULL);

char *allocptr = malloc(WIDTH*HEIGHT*2);

cl_mem buffer = clImportMemoryARM(context,
								  CL_MEM_READ_WRITE, 
								  NULL,
								  allocptr,
                                  WIDTH * HEIGHT * 2,
                                  &error );

if (error ==CL_SUCCESS)
{
	printf("sucess\n");
}
else
{
	printf("error %d\n.",error);
}

Kevin -

1) I tried toexplicitly pass CL_IMPORT_TYPE_HOST_ARM to the API using the following code. the result error code changed to -64 (i.e., #define CL_INVALID_PROPERTY                         -64).  Any hint?

char *allocptr = malloc(WIDTH*HEIGHT*2);

const cl_import_properties_arm properties[3] = {
CL_IMPORT_TYPE_ARM,  CL_IMPORT_TYPE_HOST_ARM, 
0
};


cl_mem buffer = clImportMemoryARM(context,
                                  CL_MEM_READ_WRITE, 
                                  properties, 
                                  allocptr,
                                  WIDTH*HEIGHT*2,
                                  &error);
//cl_mem buffer = clImportMemoryARM(context,CL_MEM_READ_WRITE, NULL, allocptr,WIDTH*HEIGHT*2,&error);


if (error ==CL_SUCCESS)
{
 printf("sucess\n");
}
else
{
printf("error %d\n.",error);
}

2) w.r.t. your question why we came to the conclusion that using clCreateBuffer is slower than using CPU. following are the details, hope you can shed some lights here.

a. we are using the following code to simulate multiple matrix multiplication(e.g., x times per frame at 30 fps)

for (int i = 0; i < 100; i++)
	{
		float ii = 0.12*((rand() % (10 - 1)) + 1);
		float jj = 0.4*((rand() % (10 - 1)) + 1);
		for (int j = 0; j < A.rows; j++) {
			float* data = A.ptr<float>(j);
			for (int i = 0; i < A.cols; i++) {
				data[i] = ii;
			}
		}
		for (int j = 0; j < B.rows; j++) {
			float* data = B.ptr<float>(j);
			for (int i = 0; i < B.cols; i++) {
				data[i] = jj;
			}
		}

		GPUResult = GpuFast::instance()->calcMatrix(A, B);
	}

b. the code for the GpuFast class are:

//memObjects are created as below when the class instance is created
memObjects[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * szA, NULL, NULL);
memObjects[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * szB, NULL, NULL);
memObjects[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * szC, NULL, NULL);





cv::Mat GpuFast::calcMatrix(cv::Mat a, cv::Mat b)
{
	cl_event prof_event;

	int szA = Ndim * Pdim;
	int szB = Pdim * Mdim;
	int szC = Ndim * Mdim;
	struct timeval start, end, start1, end1;
	gettimeofday(&start, NULL);

	status = clEnqueueWriteBuffer(commandQueue, memObjects[0], CL_TRUE, 0, sizeof(float) * szA, (float*)a.ptr<float>(0), 0, NULL, NULL);
	if (status) {
		printf("clEnqueueWriteBuffer  A error status=%d\n", (int)status);
	}
	status = clEnqueueWriteBuffer(commandQueue, memObjects[1], CL_TRUE, 0, sizeof(float) * szB, (float*)b.ptr<float>(0), 0, NULL, NULL);

	if (status) {
		printf("clEnqueueWriteBuffer B  error status=%d\n", (int)status);
	}

	//cl_int clnum = NWITEMS;
	status = clSetKernelArg(kernel, 0, sizeof(int), &Ndim);
	status = clSetKernelArg(kernel, 1, sizeof(int), &Mdim);
	status = clSetKernelArg(kernel, 2, sizeof(int), &Pdim);
	status = clSetKernelArg(kernel, 3, sizeof(cl_mem), &memObjects[0]);
	status = clSetKernelArg(kernel, 4, sizeof(cl_mem), &memObjects[1]);
	status = clSetKernelArg(kernel, 5, sizeof(cl_mem), &memObjects[2]);
	if (status) {
		cout << "clSetKernelArg error" << endl;
	}
	global[0] = (size_t)Ndim;
	global[1] = (size_t)Mdim;



	status = clEnqueueNDRangeKernel(
		commandQueue, kernel, 2, NULL,
		global, NULL, 0, NULL, &prof_event);
	if (status) {
		cout << "error" << endl;
	}

	gettimeofday(&end, NULL);
	printf("cpu data to gpu use time %dms\n", 1000 * (end.tv_sec - start.tv_sec) + (end.tv_usec - start.tv_usec) / 1000);

	//clFinish(commandQueue);
	clWaitForEvents(1, &prof_event);
	cl_ulong ev_start_time = (cl_ulong)0;
	cl_ulong ev_end_time = (cl_ulong)0;
	double rum_time;
	cl_int ret = clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &ev_start_time, NULL);
	ret = clGetEventProfilingInfo(prof_event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &ev_end_time, NULL);

	rum_time = ev_end_time - ev_start_time;
	printf("\nExecution time in milliseconds = %0.3f ms\n", (rum_time / 1000000.0));

	struct timeval start2, end2;
	gettimeofday(&start2, NULL);
	status = clEnqueueReadBuffer(commandQueue, memObjects[2], CL_TRUE, 0, sizeof(float)* szC, result, 0, NULL, NULL);
	if (status)
	{
		printf("clEnqueueReadBuffer error status=%d\n",(int)status);
	}
	gettimeofday(&end2, NULL);
	printf("GPU result to CPU use time %dms\n", 1000 * (end2.tv_sec - start2.tv_sec) + (end2.tv_usec - start2.tv_usec) / 1000);


	return cv::Mat(Ndim, Mdim, CV_32FC1, result);
}

c. what we realize is that clEnqueueWriteBuffer cost more or less 2ms each time. this is quite significant for delay budget, hence we are trying to look for alternative solution now.

Kevin - thanks for your valuable input. will try to optimize the code as you recommend. let me get back to you once I have some update.