//gcc ./pipes.c -o pipe -lOpenCL #include #include #include #include "CL/opencl.h" #define SIZE 1024 //------------------------------------------------------------------------------ //----------------------------------For debug----------------------------------- //------------------------------------------------------------------------------ void _trace(char const *file, long line) { printf("[%s:%ld]\n", file, line); } #define trace() _trace(__FILE__, __LINE__) const char *getErrorString(cl_int error) { switch(error) { case 0: return "CL_SUCCESS"; // run-time and JIT compiler errors case -1: return "CL_DEVICE_NOT_FOUND"; case -2: return "CL_DEVICE_NOT_AVAILABLE"; case -3: return "CL_COMPILER_NOT_AVAILABLE"; case -4: return "CL_MEM_OBJECT_ALLOCATION_FAILURE"; case -5: return "CL_OUT_OF_RESOURCES"; case -6: return "CL_OUT_OF_HOST_MEMORY"; case -7: return "CL_PROFILING_INFO_NOT_AVAILABLE"; case -8: return "CL_MEM_COPY_OVERLAP"; case -9: return "CL_IMAGE_FORMAT_MISMATCH"; case -10: return "CL_IMAGE_FORMAT_NOT_SUPPORTED"; case -11: return "CL_BUILD_PROGRAM_FAILURE"; case -12: return "CL_MAP_FAILURE"; case -13: return "CL_MISALIGNED_SUB_BUFFER_OFFSET"; case -14: return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"; case -15: return "CL_COMPILE_PROGRAM_FAILURE"; case -16: return "CL_LINKER_NOT_AVAILABLE"; case -17: return "CL_LINK_PROGRAM_FAILURE"; case -18: return "CL_DEVICE_PARTITION_FAILED"; case -19: return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE"; case -30: return "CL_INVALID_VALUE"; // compile-time errors case -31: return "CL_INVALID_DEVICE_TYPE"; case -32: return "CL_INVALID_PLATFORM"; case -33: return "CL_INVALID_DEVICE"; case -34: return "CL_INVALID_CONTEXT"; case -35: return "CL_INVALID_QUEUE_PROPERTIES"; case -36: return "CL_INVALID_COMMAND_QUEUE"; case -37: return "CL_INVALID_HOST_PTR"; case -38: return "CL_INVALID_MEM_OBJECT"; case -39: return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"; case -40: return "CL_INVALID_IMAGE_SIZE"; case -41: return "CL_INVALID_SAMPLER"; case -42: return "CL_INVALID_BINARY"; case -43: return "CL_INVALID_BUILD_OPTIONS"; case -44: return "CL_INVALID_PROGRAM"; case -45: return "CL_INVALID_PROGRAM_EXECUTABLE"; case -46: return "CL_INVALID_KERNEL_NAME"; case -47: return "CL_INVALID_KERNEL_DEFINITION"; case -48: return "CL_INVALID_KERNEL"; case -49: return "CL_INVALID_ARG_INDEX"; case -50: return "CL_INVALID_ARG_VALUE"; case -51: return "CL_INVALID_ARG_SIZE"; case -52: return "CL_INVALID_KERNEL_ARGS"; case -53: return "CL_INVALID_WORK_DIMENSION"; case -54: return "CL_INVALID_WORK_GROUP_SIZE"; case -55: return "CL_INVALID_WORK_ITEM_SIZE"; case -56: return "CL_INVALID_GLOBAL_OFFSET"; case -57: return "CL_INVALID_EVENT_WAIT_LIST"; case -58: return "CL_INVALID_EVENT"; case -59: return "CL_INVALID_OPERATION"; case -60: return "CL_INVALID_GL_OBJECT"; case -61: return "CL_INVALID_BUFFER_SIZE"; case -62: return "CL_INVALID_MIP_LEVEL"; case -63: return "CL_INVALID_GLOBAL_WORK_SIZE"; case -64: return "CL_INVALID_PROPERTY"; case -65: return "CL_INVALID_IMAGE_DESCRIPTOR"; case -66: return "CL_INVALID_COMPILER_OPTIONS"; case -67: return "CL_INVALID_LINKER_OPTIONS"; case -68: return "CL_INVALID_DEVICE_PARTITION_COUNT"; // extension errors case -1000: return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR"; case -1001: return "CL_PLATFORM_NOT_FOUND_KHR"; case -1002: return "CL_INVALID_D3D10_DEVICE_KHR"; case -1003: return "CL_INVALID_D3D10_RESOURCE_KHR"; case -1004: return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR"; case -1005: return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR"; default: return "Unknown OpenCL error"; } } void testStatus(int status) { if(status != CL_SUCCESS) { const char*real_message_errorMsg=getErrorString(status); printf("Error:%s\nError No.:%d\n", real_message_errorMsg,status); exit(EXIT_FAILURE); } } char* common_read_file(const char *path, long *length_out) { char *buffer; FILE *f; long length; f = fopen(path, "r"); fseek(f, 0, SEEK_END); length = ftell(f); fseek(f, 0, SEEK_SET); buffer = (char*)malloc(length); if (fread(buffer, 1, length, f) < (size_t)length) { return NULL; } fclose(f); if (NULL != length_out) { *length_out = length; } return buffer; } //------------------------------------------------------------------------------ //----------------------------------For debug----------------------------------- //------------------------------------------------------------------------------ int main() { FILE *fp; char fileName[] = "./kernel.cl"; char *kernel_source; size_t source_size; /* Load the source code containing the kernel*/ fp = fopen(fileName, "r"); if (!fp) { fprintf(stderr, "Failed to load kernel.\n"); exit(1); } kernel_source = (char*)malloc(SIZE); source_size = fread(kernel_source, 1, SIZE, fp); fclose(fp); int *input = (int *)malloc(sizeof(int) * SIZE); int *output = (int *)malloc(sizeof(int) * SIZE); memset(output, 0, sizeof(int) * SIZE); for (int i = 0; i != SIZE; ++i) { input[i] = i; } cl_int status; cl_platform_id platform_id[10]; cl_uint num_platforms; status = clGetPlatformIDs(10, platform_id, &num_platforms); if(status!=0) {trace();} testStatus(status); cl_device_id device; cl_uint num_devices; status = clGetDeviceIDs(platform_id[3], CL_DEVICE_TYPE_GPU, 1, &device, &num_devices); if(status!=0) {trace();} testStatus(status); cl_context context = clCreateContext(0, 1, &device, NULL, NULL, &status); if(status!=0) {trace();} testStatus(status); cl_queue_properties properties[]={CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE,0}; cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, properties, &status); if(status!=0) {trace();} testStatus(status); size_t len = strlen(kernel_source); cl_program program = clCreateProgramWithSource(context, 1, (const char **)&kernel_source, &len, &status); if(status!=0) {trace();} testStatus(status); status = clBuildProgram(program, num_devices, &device, "-cl-std=CL2.0", NULL, NULL); if (status == CL_BUILD_PROGRAM_FAILURE) { // Determine the size of the log size_t log_size; clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); // Allocate memory for the log char *log = (char *) malloc(log_size); // Get the log clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, log, NULL); // Print the log printf("%s\n", log); } if(status!=0) {trace();} testStatus(status); cl_kernel pipe_writer = clCreateKernel(program, "pipe_writer", &status); if(status!=0) {trace();} testStatus(status); cl_kernel pipe_reader = clCreateKernel(program, "pipe_reader", &status); if(status!=0) {trace();} testStatus(status); cl_mem in_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * SIZE, input, &status); if(status!=0) {trace();} testStatus(status); cl_mem out_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * SIZE, NULL, &status); if(status!=0) {trace();} testStatus(status); cl_mem pipe = clCreatePipe(context, 0, sizeof(cl_int), SIZE, NULL, &status); if(status!=0) {trace();} testStatus(status); status += clSetKernelArg(pipe_writer, 0, sizeof(cl_mem), &in_buffer); status += clSetKernelArg(pipe_writer, 1, sizeof(cl_mem), &pipe); status += clSetKernelArg(pipe_reader, 0, sizeof(cl_mem), &out_buffer); status += clSetKernelArg(pipe_reader, 1, sizeof(cl_mem), &pipe); if(status!=0) {trace();} testStatus(status); size_t size = SIZE; cl_event sync; status = clEnqueueNDRangeKernel(queue, pipe_writer, 1, NULL, &size, NULL, 0, NULL, &sync); if(status!=0) {trace();} testStatus(status); status = clEnqueueNDRangeKernel(queue, pipe_reader, 1, NULL, &size, NULL, 1, &sync, NULL); if(status!=0) {trace();} testStatus(status); status = clFinish(queue); if(status!=0) {trace();} testStatus(status); status = clEnqueueReadBuffer(queue, out_buffer, CL_TRUE, 0, sizeof(int) * SIZE, output, 0, NULL, NULL); if(status!=0) {trace();} testStatus(status); int golden = 0, result = 0; for (int i = 0; i != SIZE; ++i) { golden += input[i]; result += output[i]; printf("input:%d\t:output:%d\t\n",input[i],output[i]); } int ret = 0; if (golden != result) { printf("FAILED!"); ret = 1; } else { printf("PASSED!"); } printf("\n"); return ret; }