opencl coarse svm

设备端代码：

//using for coarse svm test
typedef struct _Element{
    __global float *internal;
    __global float *external;
    float value;
}Element;

__kernel void svmbasic(__global Element *elements, __global float *dst)
{
    int id = (int)get_global_id(0);
    float internalElement = *(elements[id].internal);
    float externalElement = *(elements[id].external);
    dst[id] = internalElement + externalElement;
}
 

主机端：

   //follow is coarse gained svm example code
    {
        typedef struct _Element
        {
            float *internal;
            float *external;
            float value;
        }Element;
        
        //create kernel
        cl_kernel kernel = clCreateKernel(program, "svmbasic", &err);
        opencl_check(err);
        
        //svm alloc buffer
        int size = 16;
        Element *inputElements = (Element*)clSVMAlloc(context, CL_MEM_READ_ONLY, sizeof(Element) * size, 0); 
        float *inputFloats = (float*)clSVMAlloc(context, CL_MEM_READ_ONLY, sizeof(float) * size, 0); 
        float *outputFloats = (float*)clSVMAlloc(context, CL_MEM_READ_ONLY, sizeof(float) * size, 0); 

        //svm mapper buffer
        err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, inputElements, sizeof(Element)*size, 0, 0, 0);
        opencl_check(err);
        err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_WRITE, inputFloats, sizeof(float)*size, 0, 0, 0);
        opencl_check(err);

        //initialize the input svm buffers
        for(int i = 0; i < size; i++)
        {
            inputElements[i].value = i;
            inputFloats[i] = i + size;
        }
        for(int i = 0; i < size; i++)
        {
            inputElements[i].internal = &inputElements[i].value; 
            inputElements[i].external = &inputFloats[i]; 
        }

        //unmap input svm buffers
        err = clEnqueueSVMUnmap(queue, inputFloats, 0, 0, 0);
        opencl_check(err);
        err = clEnqueueSVMUnmap(queue, inputElements, 0, 0, 0);
        opencl_check(err);

        //set args
        err = clSetKernelArgSVMPointer(kernel, 0, inputElements);
        opencl_check(err);
        err = clSetKernelArgSVMPointer(kernel, 1, outputFloats);
        opencl_check(err);
        //set inputFloats as executing info ,not  pass it as args
        err = clSetKernelExecInfo(kernel, CL_KERNEL_EXEC_INFO_SVM_PTRS, sizeof(inputFloats), &inputFloats);
        opencl_check(err);

        //execute kernel
        size_t globalsize = size;
        err = clEnqueueNDRangeKernel(queue, kernel, 1, 0, &globalsize, NULL, NULL, NULL, NULL); 
        opencl_check(err);
        err = clFinish(queue);
        opencl_check(err);

        //read output
        err = clEnqueueSVMMap(queue, CL_TRUE, CL_MAP_READ, outputFloats, sizeof(float) * size, 0, 0, 0);
        opencl_check(err);

        //print result
        for(int i = 0; i < size; i++)
        {
            float expectiveValue = *(inputElements[i].external) + *(inputElements[i].internal);
            printf("expect %f, real %f\n", expectiveValue, outputFloats[i]);
        }
        err = clEnqueueSVMUnmap(queue, outputFloats, 0, 0, 0);
        opencl_check(err);
        clSVMFree(context, outputFloats);
        clSVMFree(context, inputFloats);
        clSVMFree(context, inputElements);
    }