diff --git a/test_conformance/allocations/allocation_execute.cpp b/test_conformance/allocations/allocation_execute.cpp index 1762711067..692424ceb1 100644 --- a/test_conformance/allocations/allocation_execute.cpp +++ b/test_conformance/allocations/allocation_execute.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -20,7 +20,8 @@ const char *buffer_kernel_pattern = { - "__kernel void sample_test(%s __global uint *result, __global %s *array_sizes, uint per_item)\n" + "__kernel void sample_test(%s __global uint *result, __global %s " + "*array_sizes, uint per_item)\n" "{\n" "\tint tid = get_global_id(0);\n" "\tuint r = 0;\n" @@ -29,7 +30,8 @@ const char *buffer_kernel_pattern = { "%s" "\t}\n" "\tresult[tid] = r;\n" - "}\n" }; + "}\n" +}; const char *image_kernel_pattern = { "__kernel void sample_test(%s __global uint *result)\n" @@ -40,7 +42,8 @@ const char *image_kernel_pattern = { "\tint x, y;\n" "%s" "\tresult[get_global_id(0)] += color.x + color.y + color.z + color.w;\n" - "}\n" }; + "}\n" +}; const char *read_pattern = { "\tfor(y=0; y returned_results(NUM_OF_WORK_ITEMS); + std::vector returned_results(number_of_work_itmes); clEventWrapper event; cl_int event_status; // Allocate memory for the kernel source - argument_string = (char*)malloc(sizeof(char)*MAX_NUMBER_TO_ALLOCATE*64); - access_string = (char*)malloc(sizeof(char)*MAX_NUMBER_TO_ALLOCATE*(strlen(read_pattern)+10)); - kernel_string = (char*)malloc(sizeof(char)*MAX_NUMBER_TO_ALLOCATE*(strlen(read_pattern)+10+64)+1024); + argument_string = + (char *)malloc(sizeof(char) * MAX_NUMBER_TO_ALLOCATE * 64); + access_string = (char *)malloc(sizeof(char) * MAX_NUMBER_TO_ALLOCATE + * (strlen(read_pattern) + 10)); + kernel_string = (char *)malloc(sizeof(char) * MAX_NUMBER_TO_ALLOCATE + * (strlen(read_pattern) + 10 + 64) + + 1024); argument_string[0] = '\0'; access_string[0] = '\0'; kernel_string[0] = '\0'; // Zero the results. - for (i=0; i max_size) - max_size = size/sizeof(cl_uint); + if (size / sizeof(cl_uint) > max_size) + max_size = size / sizeof(cl_uint); } - if (support64) { - buffer_sizes = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_ulong)*number_of_mems_used, ulSizes, &error); + if (support64) + { + buffer_sizes = clCreateBuffer( + context, CL_MEM_COPY_HOST_PTR, + sizeof(cl_ulong) * number_of_mems_used, ulSizes, &error); } - else { - buffer_sizes = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(cl_uint)*number_of_mems_used, uiSizes, &error); + else + { + buffer_sizes = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, + sizeof(cl_uint) * number_of_mems_used, + uiSizes, &error); } test_error_abort(error, "clCreateBuffer failed"); - error = clSetKernelArg(kernel, number_of_mems_used+1, sizeof(cl_mem), &buffer_sizes); + error = clSetKernelArg(kernel, number_of_mems_used + 1, sizeof(cl_mem), + &buffer_sizes); test_error(error, "clSetKernelArg failed"); - per_item = (cl_uint)ceil((double)max_size/global_dims[0]); + per_item = (cl_uint)ceil((double)max_size / global_dims[0]); if (per_item > CL_UINT_MAX) - log_error("Size is too large for a uint parameter to the kernel. Expect invalid results.\n"); + log_error("Size is too large for a uint parameter to the kernel. " + "Expect invalid results.\n"); per_item_uint = (cl_uint)per_item; - error = clSetKernelArg(kernel, number_of_mems_used+2, sizeof(per_item_uint), &per_item_uint); + error = clSetKernelArg(kernel, number_of_mems_used + 2, + sizeof(per_item_uint), &per_item_uint); test_error(error, "clSetKernelArg failed"); } - if (ulSizes) { + if (ulSizes) + { free(ulSizes); } - if (uiSizes) { + if (uiSizes) + { free(uiSizes); } - size_t local_dims[3] = {1,1,1}; - error = get_max_common_work_group_size(context, kernel, global_dims[0], &local_dims[0]); + size_t local_dims[3] = { 1, 1, 1 }; + error = get_max_common_work_group_size(context, kernel, global_dims[0], + &local_dims[0]); test_error(error, "get_max_common_work_group_size failed"); // Execute the kernel - error = clEnqueueNDRangeKernel(*queue, kernel, 1, NULL, global_dims, local_dims, 0, NULL, &event); + error = clEnqueueNDRangeKernel(*queue, kernel, 1, NULL, global_dims, + local_dims, 0, NULL, &event); result = check_allocation_error(context, device_id, error, queue); - if (result != SUCCEEDED) { + if (result != SUCCEEDED) + { if (result == FAILED_TOO_BIG) - log_info("\t\tExecute kernel failed: %s (global dim: %ld, local dim: %ld)\n", IGetErrorString(error), global_dims[0], local_dims[0]); + log_info("\t\tExecute kernel failed: %s (global dim: %ld, local " + "dim: %ld)\n", + IGetErrorString(error), global_dims[0], local_dims[0]); else print_error(error, "clEnqueueNDRangeKernel failed"); return result; @@ -317,7 +386,8 @@ int execute_kernel(cl_context context, cl_command_queue *queue, cl_device_id dev result = check_allocation_error(context, device_id, error, queue); - if (result != SUCCEEDED) { + if (result != SUCCEEDED) + { if (result == FAILED_TOO_BIG) log_info("\t\tclFinish failed: %s.\n", IGetErrorString(error)); else @@ -326,13 +396,20 @@ int execute_kernel(cl_context context, cl_command_queue *queue, cl_device_id dev } // Verify that the event from the execution did not have an error - error = clGetEventInfo(event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(event_status), &event_status, NULL); - test_error_abort(error, "clGetEventInfo for CL_EVENT_COMMAND_EXECUTION_STATUS failed"); - if (event_status < 0) { - result = check_allocation_error(context, device_id, event_status, queue); - if (result != SUCCEEDED) { + error = clGetEventInfo(event, CL_EVENT_COMMAND_EXECUTION_STATUS, + sizeof(event_status), &event_status, NULL); + test_error_abort( + error, "clGetEventInfo for CL_EVENT_COMMAND_EXECUTION_STATUS failed"); + if (event_status < 0) + { + result = + check_allocation_error(context, device_id, event_status, queue); + if (result != SUCCEEDED) + { if (result == FAILED_TOO_BIG) - log_info("\t\tEvent returned from kernel execution indicates failure: %s.\n", IGetErrorString(event_status)); + log_info("\t\tEvent returned from kernel execution indicates " + "failure: %s.\n", + IGetErrorString(event_status)); else print_error(event_status, "clEnqueueNDRangeKernel failed"); return result; @@ -340,33 +417,46 @@ int execute_kernel(cl_context context, cl_command_queue *queue, cl_device_id dev } // If we are not verifying the checksum return here - if (!verify_checksum) { - log_info("Note: Allocations were not initialized so kernel execution can not verify correct results.\n"); + if (!verify_checksum) + { + log_info("Note: Allocations were not initialized so kernel execution " + "can not verify correct results.\n"); return SUCCEEDED; } // Verify the checksum. // Read back the result error = clEnqueueReadBuffer(*queue, result_mem, CL_TRUE, 0, - sizeof(cl_uint) * NUM_OF_WORK_ITEMS, + sizeof(cl_uint) * number_of_work_itmes, returned_results.data(), 0, NULL, NULL); test_error_abort(error, "clEnqueueReadBuffer failed"); final_result = 0; - if (test == BUFFER || test == IMAGE_READ || test == BUFFER_NON_BLOCKING || test == IMAGE_READ_NON_BLOCKING) { - // For buffers or read images we are just looking at the sum of what each thread summed up - for (i=0; i size) - size_to_use = size; - - data = (cl_uint*)malloc(size_to_use); - if (data == NULL) { - log_error("Failed to malloc host buffer for writing into buffer.\n"); - return FAILED_ABORT; - } - for (i=0; i size) size_to_use = size; + + data = (cl_uint *)malloc(size_to_use); + if (data == NULL) + { + log_error("Failed to malloc host buffer for writing into buffer.\n"); + return FAILED_ABORT; } - } - - // Deal with any leftover bits - if (i < size) { - // Put values in the data, and keep a checksum as we go along. - for (j=0; j<(size-i)/sizeof(cl_uint); j++) { - data[j] = (cl_uint)genrand_int32(d); - checksum_delta += data[j]; + for (i = 0; i < size - size_to_use; i += size_to_use) + { + // Put values in the data, and keep a checksum as we go along. + for (j = 0; j < size_to_use / sizeof(cl_uint); j++) + { + data[j] = genrand_int32(d); + checksum_delta += data[j]; + } + if (blocking_write) + { + error = clEnqueueWriteBuffer(*queue, mem, CL_TRUE, i, size_to_use, + data, 0, NULL, NULL); + result = check_allocation_error(context, device_id, error, queue); + + if (result == FAILED_ABORT) + { + print_error(error, "clEnqueueWriteBuffer failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + free(data); + clReleaseMemObject(mem); + return result; + } + } + else + { + error = clEnqueueWriteBuffer(*queue, mem, CL_FALSE, i, size_to_use, + data, 0, NULL, &event); + result = check_allocation_error(context, device_id, error, queue); + + if (result == FAILED_ABORT) + { + print_error(error, "clEnqueueWriteBuffer failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + free(data); + clReleaseMemObject(mem); + return result; + } + + error = clWaitForEvents(1, &event); + result = check_allocation_error(context, device_id, error, queue, + &event); + + if (result == FAILED_ABORT) + { + print_error(error, "clWaitForEvents failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + clReleaseEvent(event); + free(data); + clReleaseMemObject(mem); + return result; + } + + clReleaseEvent(event); + } } - if (blocking_write) { - error = clEnqueueWriteBuffer(*queue, mem, CL_TRUE, i, size-i, data, 0, NULL, NULL); - result = check_allocation_error(context, device_id, error, queue); - - if (result == FAILED_ABORT) { - print_error(error, "clEnqueueWriteBuffer failed."); - } - - if (result != SUCCEEDED) { - clFinish(*queue); - clReleaseMemObject(mem); - free(data); - return result; - } - } else { - error = clEnqueueWriteBuffer(*queue, mem, CL_FALSE, i, size-i, data, 0, NULL, &event); - result = check_allocation_error(context, device_id, error, queue); - - if (result == FAILED_ABORT) { - print_error(error, "clEnqueueWriteBuffer failed."); - } - - if (result != SUCCEEDED) { - clFinish(*queue); - clReleaseMemObject(mem); - free(data); - return result; - } - - error = clWaitForEvents(1, &event); - result = check_allocation_error(context, device_id, error, queue, &event); - - if (result == FAILED_ABORT) { - print_error(error, "clWaitForEvents failed."); - } - - if (result != SUCCEEDED) { - clFinish(*queue); - clReleaseEvent(event); - free(data); - clReleaseMemObject(mem); - return result; - } - - clReleaseEvent(event); + // Deal with any leftover bits + if (i < size) + { + // Put values in the data, and keep a checksum as we go along. + for (j = 0; j < (size - i) / sizeof(cl_uint); j++) + { + data[j] = (cl_uint)genrand_int32(d); + checksum_delta += data[j]; + } + + if (blocking_write) + { + error = clEnqueueWriteBuffer(*queue, mem, CL_TRUE, i, size - i, + data, 0, NULL, NULL); + result = check_allocation_error(context, device_id, error, queue); + + if (result == FAILED_ABORT) + { + print_error(error, "clEnqueueWriteBuffer failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + clReleaseMemObject(mem); + free(data); + return result; + } + } + else + { + error = clEnqueueWriteBuffer(*queue, mem, CL_FALSE, i, size - i, + data, 0, NULL, &event); + result = check_allocation_error(context, device_id, error, queue); + + if (result == FAILED_ABORT) + { + print_error(error, "clEnqueueWriteBuffer failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + clReleaseMemObject(mem); + free(data); + return result; + } + + error = clWaitForEvents(1, &event); + result = check_allocation_error(context, device_id, error, queue, + &event); + + if (result == FAILED_ABORT) + { + print_error(error, "clWaitForEvents failed."); + } + + if (result != SUCCEEDED) + { + clFinish(*queue); + clReleaseEvent(event); + free(data); + clReleaseMemObject(mem); + return result; + } + + clReleaseEvent(event); + } } - } - free(data); - // Only update the checksum if this succeeded. - checksum += checksum_delta; - return SUCCEEDED; + free(data); + // Only update the checksum if this succeeded. + checksum += checksum_delta; + return SUCCEEDED; } -int fill_image_with_data(cl_context context, cl_device_id device_id, cl_command_queue *queue, cl_mem mem, size_t width, size_t height, MTdata d, cl_bool blocking_write) { - size_t origin[3], region[3], j; - int error, result; - cl_uint *data; - cl_uint checksum_delta = 0; - cl_event event; - - size_t image_lines_to_use; - image_lines_to_use = IMAGE_LINES; - if (image_lines_to_use > height) - image_lines_to_use = height; - - data = (cl_uint*)malloc(width*4*sizeof(cl_uint)*image_lines_to_use); - if (data == NULL) { - log_error("Failed to malloc host buffer for writing into image.\n"); - return FAILED_ABORT; - } - origin[0] = 0; - origin[1] = 0; - origin[2] = 0; - region[0] = width; - region[1] = image_lines_to_use; - region[2] = 1; - for (origin[1] = 0; origin[1] < height - image_lines_to_use; origin[1] += image_lines_to_use) { - // Put values in the data, and keep a checksum as we go along. - for (j=0; j height) image_lines_to_use = height; + + data = (cl_uint *)malloc(width * 4 * sizeof(cl_uint) * image_lines_to_use); + if (data == NULL) + { + log_error("Failed to malloc host buffer for writing into image.\n"); + return FAILED_ABORT; } - - if (blocking_write) { - error = clEnqueueWriteImage(*queue, mem, CL_TRUE, origin, region, 0, 0, data, 0, NULL, NULL); - result = check_allocation_error(context, device_id, error, queue); - - if (result == FAILED_ABORT) { - print_error(error, "clEnqueueWriteImage failed."); - } - - if (result != SUCCEEDED) { - clFinish(*queue); - clReleaseMemObject(mem); - free(data); - return result; - } - result = clFinish(*queue); - if (result != SUCCEEDED) - { - print_error(error, - "clFinish failed after successful enqueuing filling " - "buffer with data."); - return result; - } - } else { - error = clEnqueueWriteImage(*queue, mem, CL_FALSE, origin, region, 0, 0, data, 0, NULL, &event); - result = check_allocation_error(context, device_id, error, queue); - - if (result == FAILED_ABORT) { - print_error(error, "clEnqueueWriteImage failed."); - } - - if (result != SUCCEEDED) { - clFinish(*queue); - clReleaseMemObject(mem); - free(data); - return result; - } - - error = clWaitForEvents(1, &event); - result = check_allocation_error(context, device_id, error, queue, &event); - - if (result == FAILED_ABORT) { - print_error(error, "clWaitForEvents failed."); - } - - if (result != SUCCEEDED) { - clReleaseEvent(event); - free(data); - clReleaseMemObject(mem); - return result; - } - - clReleaseEvent(event); - } - } - - // Deal with any leftover bits - if (origin[1] < height) { - // Put values in the data, and keep a checksum as we go along. - for (j=0; j max_pixels) + { + if (NULL != max_size) + { + *max_size = max_width * max_height * sizeof(cl_uint) * 4; + } + return FAILED_TOO_BIG; + } - if (num_pixels > max_pixels) { - if(NULL != max_size) { - *max_size = max_width * max_height * sizeof(cl_uint) * 4; + // We want a close-to-square aspect ratio. + // Note that this implicitly assumes that max width >= max height + found_width = (int)sqrt((double)num_pixels); + if (found_width > max_width) + { + found_width = max_width; } - return FAILED_TOO_BIG; - } - - // We want a close-to-square aspect ratio. - // Note that this implicitly assumes that max width >= max height - found_width = (int)sqrt( (double) num_pixels ); - if( found_width > max_width ) { - found_width = max_width; - } - if (found_width == 0) - found_width = 1; - - found_height = (size_t)num_pixels/found_width; - if (found_height > max_height) { - found_height = max_height; - } - if (found_height == 0) - found_height = 1; - - *width = found_width; - *height = found_height; - - if(NULL != max_size) { - *max_size = found_width * found_height * sizeof(cl_uint) * 4; - } - - return SUCCEEDED; + if (found_width == 0) found_width = 1; + + found_height = (size_t)num_pixels / found_width; + if (found_height > max_height) + { + found_height = max_height; + } + if (found_height == 0) found_height = 1; + + *width = found_width; + *height = found_height; + + if (NULL != max_size) + { + *max_size = found_width * found_height * sizeof(cl_uint) * 4; + } + + return SUCCEEDED; } -int allocate_image2d_read(cl_context context, cl_command_queue *queue, cl_device_id device_id, cl_mem *mem, size_t size_to_allocate, cl_bool blocking_write) { - size_t width, height; - int error; +int allocate_image2d_read(cl_context context, cl_command_queue *queue, + cl_device_id device_id, cl_mem *mem, + size_t size_to_allocate, cl_bool blocking_write) +{ + size_t width, height; + int error; - error = find_good_image_size(device_id, size_to_allocate, &width, &height, NULL); - if (error != SUCCEEDED) - return error; + error = find_good_image_size(device_id, size_to_allocate, &width, &height, + NULL); + if (error != SUCCEEDED) return error; - log_info("\t\tAttempting to allocate a %gMB read-only image (%d x %d) and fill with %s writes.\n", - (size_to_allocate/(1024.0*1024.0)), (int)width, (int)height, (blocking_write ? "blocking" : "non-blocking")); - *mem = create_image_2d(context, CL_MEM_READ_ONLY, &image_format, width, height, 0, NULL, &error); + log_info("\t\tAttempting to allocate a %gMB read-only image (%d x %d) and " + "fill with %s writes.\n", + (size_to_allocate / (1024.0 * 1024.0)), (int)width, (int)height, + (blocking_write ? "blocking" : "non-blocking")); + *mem = create_image_2d(context, CL_MEM_READ_ONLY, &image_format, width, + height, 0, NULL, &error); - return check_allocation_error(context, device_id, error, queue); + return check_allocation_error(context, device_id, error, queue); } -int allocate_image2d_write(cl_context context, cl_command_queue *queue, cl_device_id device_id, cl_mem *mem, size_t size_to_allocate, cl_bool blocking_write) { - size_t width, height; - int error; +int allocate_image2d_write(cl_context context, cl_command_queue *queue, + cl_device_id device_id, cl_mem *mem, + size_t size_to_allocate, cl_bool blocking_write) +{ + size_t width, height; + int error; - error = find_good_image_size(device_id, size_to_allocate, &width, &height, NULL); - if (error != SUCCEEDED) - return error; + error = find_good_image_size(device_id, size_to_allocate, &width, &height, + NULL); + if (error != SUCCEEDED) return error; - //log_info("\t\tAttempting to allocate a %gMB write-only image (%d x %d) and fill with %s writes.\n", - //(size_to_allocate/(1024.0*1024.0)), (int)width, (int)height, (blocking_write ? "blocking" : "non-blocking")); - *mem = create_image_2d(context, CL_MEM_WRITE_ONLY, &image_format, width, height, 0, NULL, &error); + // log_info("\t\tAttempting to allocate a %gMB write-only image (%d x %d) + // and fill with %s writes.\n", (size_to_allocate/(1024.0*1024.0)), + //(int)width, (int)height, (blocking_write ? "blocking" : "non-blocking")); + *mem = create_image_2d(context, CL_MEM_WRITE_ONLY, &image_format, width, + height, 0, NULL, &error); - return check_allocation_error(context, device_id, error, queue); + return check_allocation_error(context, device_id, error, queue); } -int do_allocation(cl_context context, cl_command_queue *queue, cl_device_id device_id, size_t size_to_allocate, int type, cl_mem *mem) { - if (type == BUFFER) return allocate_buffer(context, queue, device_id, mem, size_to_allocate, true); - if (type == IMAGE_READ) return allocate_image2d_read(context, queue, device_id, mem, size_to_allocate, true); - if (type == IMAGE_WRITE) return allocate_image2d_write(context, queue, device_id, mem, size_to_allocate, true); - if (type == BUFFER_NON_BLOCKING) return allocate_buffer(context, queue, device_id, mem, size_to_allocate, false); - if (type == IMAGE_READ_NON_BLOCKING) return allocate_image2d_read(context, queue, device_id, mem, size_to_allocate, false); - if (type == IMAGE_WRITE_NON_BLOCKING) return allocate_image2d_write(context, queue, device_id, mem, size_to_allocate, false); +int do_allocation(cl_context context, cl_command_queue *queue, + cl_device_id device_id, size_t size_to_allocate, int type, + cl_mem *mem) +{ + if (type == BUFFER) + return allocate_buffer(context, queue, device_id, mem, size_to_allocate, + true); + if (type == IMAGE_READ) + return allocate_image2d_read(context, queue, device_id, mem, + size_to_allocate, true); + if (type == IMAGE_WRITE) + return allocate_image2d_write(context, queue, device_id, mem, + size_to_allocate, true); + if (type == BUFFER_NON_BLOCKING) + return allocate_buffer(context, queue, device_id, mem, size_to_allocate, + false); + if (type == IMAGE_READ_NON_BLOCKING) + return allocate_image2d_read(context, queue, device_id, mem, + size_to_allocate, false); + if (type == IMAGE_WRITE_NON_BLOCKING) + return allocate_image2d_write(context, queue, device_id, mem, + size_to_allocate, false); log_error("Invalid allocation type: %d\n", type); - return FAILED_ABORT; + return FAILED_ABORT; } -int allocate_size(cl_context context, cl_command_queue *queue, cl_device_id device_id, int multiple_allocations, size_t size_to_allocate, - int type, cl_mem mems[], int *number_of_mems, size_t *final_size, int force_fill, MTdata d) { +int allocate_size(cl_context context, cl_command_queue *queue, + cl_device_id device_id, int multiple_allocations, + size_t size_to_allocate, int type, cl_mem mems[], + int *number_of_mems, size_t *final_size, int force_fill, + MTdata d) +{ cl_ulong max_individual_allocation_size, global_mem_size; - int error, result; - size_t amount_allocated; - size_t reduction_amount; - int current_allocation; - size_t allocation_this_time, actual_allocation; - - // Set the number of mems used to 0 so if we fail to create even a single one we don't end up returning a garbage value - *number_of_mems = 0; - - error = clGetDeviceInfo(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(max_individual_allocation_size), &max_individual_allocation_size, NULL); - test_error_abort( error, "clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE"); - error = clGetDeviceInfo(device_id, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(global_mem_size), &global_mem_size, NULL); - test_error_abort( error, "clGetDeviceInfo failed for CL_DEVICE_GLOBAL_MEM_SIZE"); - - if (global_mem_size > (cl_ulong)SIZE_MAX) { - global_mem_size = (cl_ulong)SIZE_MAX; - } - -// log_info("Device reports CL_DEVICE_MAX_MEM_ALLOC_SIZE=%llu bytes (%gMB), CL_DEVICE_GLOBAL_MEM_SIZE=%llu bytes (%gMB).\n", -// max_individual_allocation_size, toMB(max_individual_allocation_size), -// global_mem_size, toMB(global_mem_size)); - - if (size_to_allocate > global_mem_size) { - log_error("Can not allocate more than the global memory size.\n"); - return FAILED_ABORT; - } - - amount_allocated = 0; - current_allocation = 0; - - // If allocating for images, reduce the maximum allocation size to the maximum image size. - // If we don't do this, then the value of CL_DEVICE_MAX_MEM_ALLOC_SIZE / 4 can be higher - // than the maximum image size on systems with 16GB or RAM or more. In this case, we - // succeed in allocating an image but its size is less than CL_DEVICE_MAX_MEM_ALLOC_SIZE / 4 - // (min_allocation_allowed) and thus we fail the allocation below. - if(type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING || type == IMAGE_WRITE || type == IMAGE_WRITE_NON_BLOCKING) { - size_t width; - size_t height; - size_t max_size; - error = find_good_image_size(device_id, size_to_allocate, &width, &height, &max_size); - if (!(error == SUCCEEDED || error == FAILED_TOO_BIG)) - return error; - if(max_size < max_individual_allocation_size) - max_individual_allocation_size = max_size; - } - - reduction_amount = (size_t)max_individual_allocation_size/16; - - if (type == BUFFER || type == BUFFER_NON_BLOCKING) log_info("\tAttempting to allocate a buffer of size %gMB.\n", toMB(size_to_allocate)); - else if (type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING) log_info("\tAttempting to allocate a read-only image of size %gMB.\n", toMB(size_to_allocate)); - else if (type == IMAGE_WRITE || type == IMAGE_WRITE_NON_BLOCKING) log_info("\tAttempting to allocate a write-only image of size %gMB.\n", toMB(size_to_allocate)); - -// log_info("\t\t(Reduction size is %gMB per iteration, minimum allowable individual allocation size is %gMB.)\n", -// toMB(reduction_amount), toMB(min_allocation_allowed)); -// if (force_fill && type != IMAGE_WRITE && type != IMAGE_WRITE_NON_BLOCKING) log_info("\t\t(Allocations will be filled with random data for checksum calculation.)\n"); - - // If we are only doing a single allocation, only allow 1 - int max_to_allocate = multiple_allocations ? MAX_NUMBER_TO_ALLOCATE : 1; - - // Make sure that the maximum number of images allocated is constrained by the - // maximum that may be passed to a kernel - if (type != BUFFER && type != BUFFER_NON_BLOCKING) { - cl_device_info param_name = (type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING) ? - CL_DEVICE_MAX_READ_IMAGE_ARGS : CL_DEVICE_MAX_WRITE_IMAGE_ARGS; - - cl_uint max_image_args; - error = clGetDeviceInfo(device_id, param_name, sizeof(max_image_args), &max_image_args, NULL); - test_error( error, "clGetDeviceInfo failed for CL_DEVICE_MAX IMAGE_ARGS"); - - if ((int)max_image_args < max_to_allocate) { - log_info("\t\tMaximum number of images per kernel limited to %d\n",(int)max_image_args); - max_to_allocate = max_image_args; + int error, result; + size_t amount_allocated; + size_t reduction_amount; + int current_allocation; + size_t allocation_this_time, actual_allocation; + + // Set the number of mems used to 0 so if we fail to create even a single + // one we don't end up returning a garbage value + *number_of_mems = 0; + + error = clGetDeviceInfo(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(max_individual_allocation_size), + &max_individual_allocation_size, NULL); + test_error_abort(error, + "clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE"); + error = clGetDeviceInfo(device_id, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(global_mem_size), &global_mem_size, NULL); + test_error_abort(error, + "clGetDeviceInfo failed for CL_DEVICE_GLOBAL_MEM_SIZE"); + + if (global_mem_size > (cl_ulong)SIZE_MAX) + { + global_mem_size = (cl_ulong)SIZE_MAX; } - } + // log_info("Device reports CL_DEVICE_MAX_MEM_ALLOC_SIZE=%llu bytes (%gMB), + // CL_DEVICE_GLOBAL_MEM_SIZE=%llu bytes (%gMB).\n", + // max_individual_allocation_size, + // toMB(max_individual_allocation_size), global_mem_size, + // toMB(global_mem_size)); - // Try to allocate the requested amount. - while (amount_allocated != size_to_allocate && current_allocation < max_to_allocate) { - - // Determine how much more is needed - allocation_this_time = size_to_allocate - amount_allocated; + if (size_to_allocate > global_mem_size) + { + log_error("Can not allocate more than the global memory size.\n"); + return FAILED_ABORT; + } - // Bound by the individual allocation size - if (allocation_this_time > max_individual_allocation_size) - allocation_this_time = (size_t)max_individual_allocation_size; + amount_allocated = 0; + current_allocation = 0; + + // If allocating for images, reduce the maximum allocation size to the + // maximum image size. If we don't do this, then the value of + // CL_DEVICE_MAX_MEM_ALLOC_SIZE / 4 can be higher than the maximum image + // size on systems with 16GB or RAM or more. In this case, we succeed in + // allocating an image but its size is less than + // CL_DEVICE_MAX_MEM_ALLOC_SIZE / 4 (min_allocation_allowed) and thus we + // fail the allocation below. + if (type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING + || type == IMAGE_WRITE || type == IMAGE_WRITE_NON_BLOCKING) + { + size_t width; + size_t height; + size_t max_size; + error = find_good_image_size(device_id, size_to_allocate, &width, + &height, &max_size); + if (!(error == SUCCEEDED || error == FAILED_TOO_BIG)) return error; + if (max_size < max_individual_allocation_size) + max_individual_allocation_size = max_size; + } - // Allocate the largest object possible - result = FAILED_TOO_BIG; - //log_info("\t\tTrying sub-allocation %d at size %gMB.\n", current_allocation, toMB(allocation_this_time)); - while (result == FAILED_TOO_BIG && allocation_this_time != 0) { + reduction_amount = (size_t)max_individual_allocation_size / 16; + + if (type == BUFFER || type == BUFFER_NON_BLOCKING) + log_info("\tAttempting to allocate a buffer of size %gMB.\n", + toMB(size_to_allocate)); + else if (type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING) + log_info("\tAttempting to allocate a read-only image of size %gMB.\n", + toMB(size_to_allocate)); + else if (type == IMAGE_WRITE || type == IMAGE_WRITE_NON_BLOCKING) + log_info("\tAttempting to allocate a write-only image of size %gMB.\n", + toMB(size_to_allocate)); + + // log_info("\t\t(Reduction size is %gMB per iteration, minimum allowable + // individual allocation size is %gMB.)\n", + // toMB(reduction_amount), toMB(min_allocation_allowed)); + // if (force_fill && type != IMAGE_WRITE && type != + // IMAGE_WRITE_NON_BLOCKING) log_info("\t\t(Allocations will be filled with + // random data for checksum calculation.)\n"); + + // If we are only doing a single allocation, only allow 1 + int max_to_allocate = multiple_allocations ? MAX_NUMBER_TO_ALLOCATE : 1; + + // Make sure that the maximum number of images allocated is constrained by + // the maximum that may be passed to a kernel + if (type != BUFFER && type != BUFFER_NON_BLOCKING) + { + cl_device_info param_name = + (type == IMAGE_READ || type == IMAGE_READ_NON_BLOCKING) + ? CL_DEVICE_MAX_READ_IMAGE_ARGS + : CL_DEVICE_MAX_WRITE_IMAGE_ARGS; + + cl_uint max_image_args; + error = clGetDeviceInfo(device_id, param_name, sizeof(max_image_args), + &max_image_args, NULL); + test_error(error, + "clGetDeviceInfo failed for CL_DEVICE_MAX IMAGE_ARGS"); + + if ((int)max_image_args < max_to_allocate) + { + log_info("\t\tMaximum number of images per kernel limited to %d\n", + (int)max_image_args); + max_to_allocate = max_image_args; + } + } - // Create the object - result = do_allocation(context, queue, device_id, allocation_this_time, type, &mems[current_allocation]); - if (result == SUCCEEDED) { - // Allocation succeeded, another memory object was added to the array - *number_of_mems = (current_allocation+1); - // Verify the size is correct to within 1MB. - actual_allocation = get_actual_allocation_size(mems[current_allocation]); - if (fabs((double)allocation_this_time - (double)actual_allocation) > 1024.0*1024.0) { - log_error("Allocation not of expected size. Expected %gMB, got %gMB.\n", toMB(allocation_this_time), toMB( actual_allocation)); - return FAILED_ABORT; + // Try to allocate the requested amount. + while (amount_allocated != size_to_allocate + && current_allocation < max_to_allocate) + { + + // Determine how much more is needed + allocation_this_time = size_to_allocate - amount_allocated; + + // Bound by the individual allocation size + if (allocation_this_time > max_individual_allocation_size) + allocation_this_time = (size_t)max_individual_allocation_size; + + // Allocate the largest object possible + result = FAILED_TOO_BIG; + // log_info("\t\tTrying sub-allocation %d at size %gMB.\n", + // current_allocation, toMB(allocation_this_time)); + while (result == FAILED_TOO_BIG && allocation_this_time != 0) + { + + // Create the object + result = + do_allocation(context, queue, device_id, allocation_this_time, + type, &mems[current_allocation]); + if (result == SUCCEEDED) + { + // Allocation succeeded, another memory object was added to the + // array + *number_of_mems = (current_allocation + 1); + + // Verify the size is correct to within 1MB. + actual_allocation = + get_actual_allocation_size(mems[current_allocation]); + if (fabs((double)allocation_this_time + - (double)actual_allocation) + > 1024.0 * 1024.0) + { + log_error("Allocation not of expected size. Expected %gMB, " + "got %gMB.\n", + toMB(allocation_this_time), + toMB(actual_allocation)); + return FAILED_ABORT; + } + + // If we are filling the allocation for verification do so + if (force_fill) + { + // log_info("\t\t\tWriting random values to object and + // calculating checksum.\n"); + cl_bool blocking_write = true; + if (type == BUFFER_NON_BLOCKING + || type == IMAGE_READ_NON_BLOCKING + || type == IMAGE_WRITE_NON_BLOCKING) + { + blocking_write = false; + } + result = fill_mem_with_data(context, device_id, queue, + mems[current_allocation], d, + blocking_write); + } + } + + // If creation failed, try to create a smaller object + if (result == FAILED_TOO_BIG) + { + // log_info("\t\t\tAllocation %d failed at size %gMB. Trying + // smaller.\n", current_allocation, toMB(allocation_this_time)); + if (allocation_this_time > reduction_amount) + allocation_this_time -= reduction_amount; + else if (reduction_amount > 1) + { + reduction_amount /= 2; + } + else + { + allocation_this_time = 0; + } + } } - // If we are filling the allocation for verification do so - if (force_fill) { - //log_info("\t\t\tWriting random values to object and calculating checksum.\n"); - cl_bool blocking_write = true; - if (type == BUFFER_NON_BLOCKING || type == IMAGE_READ_NON_BLOCKING || type == IMAGE_WRITE_NON_BLOCKING) { - blocking_write = false; - } - result = fill_mem_with_data(context, device_id, queue, mems[current_allocation], d, blocking_write); - } - } - - // If creation failed, try to create a smaller object - if (result == FAILED_TOO_BIG) { - //log_info("\t\t\tAllocation %d failed at size %gMB. Trying smaller.\n", current_allocation, toMB(allocation_this_time)); - if (allocation_this_time > reduction_amount) - allocation_this_time -= reduction_amount; - else if (reduction_amount > 1) { - reduction_amount /= 2; - } - else { - allocation_this_time = 0; + if (result == FAILED_ABORT) + { + log_error("\t\tAllocation failed.\n"); + return FAILED_ABORT; } - } - } + if (!allocation_this_time) + { + log_info("\t\tFailed to allocate %gMB across several objects.\n", + toMB(size_to_allocate)); + return FAILED_TOO_BIG; + } - if (result == FAILED_ABORT) { - log_error("\t\tAllocation failed.\n"); - return FAILED_ABORT; - } + // Otherwise we succeeded + if (result != SUCCEEDED) + { + log_error("Test logic error."); + exit(-1); + } + amount_allocated += allocation_this_time; - if (!allocation_this_time) { - log_info("\t\tFailed to allocate %gMB across several objects.\n", toMB(size_to_allocate)); - return FAILED_TOO_BIG; - } + *final_size = amount_allocated; - // Otherwise we succeeded - if (result != SUCCEEDED) { - log_error("Test logic error."); - exit(-1); + current_allocation++; } - amount_allocated += allocation_this_time; - - *final_size = amount_allocated; - - current_allocation++; - } - log_info("\t\tSucceeded in allocating %gMB using %d memory objects.\n", toMB(amount_allocated), current_allocation); - return SUCCEEDED; + log_info("\t\tSucceeded in allocating %gMB using %d memory objects.\n", + toMB(amount_allocated), current_allocation); + return SUCCEEDED; } diff --git a/test_conformance/allocations/allocation_functions.h b/test_conformance/allocations/allocation_functions.h index 939a993bba..d93a09209c 100644 --- a/test_conformance/allocations/allocation_functions.h +++ b/test_conformance/allocations/allocation_functions.h @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -16,9 +16,20 @@ #include "testBase.h" #include "allocation_utils.h" -int do_allocation(cl_context context, cl_command_queue *queue, cl_device_id device_id, size_t size_to_allocate, int type, cl_mem *mem); -int allocate_buffer(cl_context context, cl_command_queue *queue, cl_device_id device_id, cl_mem *mem, size_t size_to_allocate); -int allocate_image2d_read(cl_context context, cl_command_queue *queue, cl_device_id device_id, cl_mem *mem, size_t size_to_allocate); -int allocate_image2d_write(cl_context context, cl_command_queue *queue, cl_device_id device_id, cl_mem *mem, size_t size_to_allocate); -int allocate_size(cl_context context, cl_command_queue *queue, cl_device_id device_id, int multiple_allocations, size_t size_to_allocate, - int type, cl_mem mems[], int *number_of_mems, size_t *final_size, int force_fill, MTdata d); +int do_allocation(cl_context context, cl_command_queue *queue, + cl_device_id device_id, size_t size_to_allocate, int type, + cl_mem *mem); +int allocate_buffer(cl_context context, cl_command_queue *queue, + cl_device_id device_id, cl_mem *mem, + size_t size_to_allocate); +int allocate_image2d_read(cl_context context, cl_command_queue *queue, + cl_device_id device_id, cl_mem *mem, + size_t size_to_allocate); +int allocate_image2d_write(cl_context context, cl_command_queue *queue, + cl_device_id device_id, cl_mem *mem, + size_t size_to_allocate); +int allocate_size(cl_context context, cl_command_queue *queue, + cl_device_id device_id, int multiple_allocations, + size_t size_to_allocate, int type, cl_mem mems[], + int *number_of_mems, size_t *final_size, int force_fill, + MTdata d); diff --git a/test_conformance/allocations/allocation_utils.cpp b/test_conformance/allocations/allocation_utils.cpp index 7d6520b0e4..95575750f6 100644 --- a/test_conformance/allocations/allocation_utils.cpp +++ b/test_conformance/allocations/allocation_utils.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -15,90 +15,116 @@ // #include "allocation_utils.h" -cl_command_queue reset_queue(cl_context context, cl_device_id device_id, cl_command_queue *queue, int *error) +cl_command_queue reset_queue(cl_context context, cl_device_id device_id, + cl_command_queue *queue, int *error) { - log_info("Invalid command queue. Releasing and recreating the command queue.\n"); - clReleaseCommandQueue(*queue); + log_info( + "Invalid command queue. Releasing and recreating the command queue.\n"); + clReleaseCommandQueue(*queue); *queue = clCreateCommandQueue(context, device_id, 0, error); - return *queue; + return *queue; } -int check_allocation_error(cl_context context, cl_device_id device_id, int error, cl_command_queue *queue, cl_event *event) { - //log_info("check_allocation_error context=%p device_id=%p error=%d *queue=%p\n", context, device_id, error, *queue); - if (error == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST && event != 0) - { - // check for errors from clWaitForEvents (e.g after clEnqueueWriteBuffer) - cl_int eventError; - error = clGetEventInfo(*event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(error), &eventError, 0); - if (CL_SUCCESS != error) +int check_allocation_error(cl_context context, cl_device_id device_id, + int error, cl_command_queue *queue, cl_event *event) +{ + // log_info("check_allocation_error context=%p device_id=%p error=%d + // *queue=%p\n", context, device_id, error, *queue); + if (error == CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST && event != 0) + { + // check for errors from clWaitForEvents (e.g after + // clEnqueueWriteBuffer) + cl_int eventError; + error = clGetEventInfo(*event, CL_EVENT_COMMAND_EXECUTION_STATUS, + sizeof(error), &eventError, 0); + if (CL_SUCCESS != error) + { + log_error("Failed to get event execution status: %s\n", + IGetErrorString(error)); + return FAILED_ABORT; + } + if (eventError >= 0) + { + log_error("Non-negative event execution status after " + "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: %s\n", + IGetErrorString(error)); + return FAILED_ABORT; + } + error = eventError; + } + if ((error == CL_MEM_OBJECT_ALLOCATION_FAILURE) + || (error == CL_OUT_OF_RESOURCES) || (error == CL_OUT_OF_HOST_MEMORY) + || (error == CL_INVALID_IMAGE_SIZE)) { - log_error("Failed to get event execution status: %s\n", IGetErrorString(error)); - return FAILED_ABORT; + return FAILED_TOO_BIG; } - if (eventError >= 0) + else if (error == CL_INVALID_COMMAND_QUEUE) { - log_error("Non-negative event execution status after CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: %s\n", IGetErrorString(error)); - return FAILED_ABORT; + *queue = reset_queue(context, device_id, queue, &error); + if (CL_SUCCESS != error) + { + log_error( + "Failed to reset command queue after corrupted queue: %s\n", + IGetErrorString(error)); + return FAILED_ABORT; + } + // Try again with smaller resources. + return FAILED_TOO_BIG; } - error = eventError; - } - if ((error == CL_MEM_OBJECT_ALLOCATION_FAILURE ) || (error == CL_OUT_OF_RESOURCES ) || (error == CL_OUT_OF_HOST_MEMORY) || (error == CL_INVALID_IMAGE_SIZE)) { - return FAILED_TOO_BIG; - } else if (error == CL_INVALID_COMMAND_QUEUE) { - *queue = reset_queue(context, device_id, queue, &error); - if (CL_SUCCESS != error) + else if (error != CL_SUCCESS) { - log_error("Failed to reset command queue after corrupted queue: %s\n", IGetErrorString(error)); - return FAILED_ABORT; + log_error("Allocation failed with %s.\n", IGetErrorString(error)); + return FAILED_ABORT; } - // Try again with smaller resources. - return FAILED_TOO_BIG; - } else if (error != CL_SUCCESS) { - log_error("Allocation failed with %s.\n", IGetErrorString(error)); - return FAILED_ABORT; - } - return SUCCEEDED; + return SUCCEEDED; } -double toMB(cl_ulong size_in) { - return (double)size_in/(1024.0*1024.0); -} - -size_t get_actual_allocation_size(cl_mem mem) { - int error; - cl_mem_object_type type; - size_t size, width, height; +double toMB(cl_ulong size_in) { return (double)size_in / (1024.0 * 1024.0); } - error = clGetMemObjectInfo(mem, CL_MEM_TYPE, sizeof(type), &type, NULL); - if (error) { - print_error(error, "clGetMemObjectInfo failed for CL_MEM_TYPE."); - return 0; - } +size_t get_actual_allocation_size(cl_mem mem) +{ + int error; + cl_mem_object_type type; + size_t size, width, height; - if (type == CL_MEM_OBJECT_BUFFER) { - error = clGetMemObjectInfo(mem, CL_MEM_SIZE, sizeof(size), &size, NULL); - if (error) { - print_error(error, "clGetMemObjectInfo failed for CL_MEM_SIZE."); - return 0; + error = clGetMemObjectInfo(mem, CL_MEM_TYPE, sizeof(type), &type, NULL); + if (error) + { + print_error(error, "clGetMemObjectInfo failed for CL_MEM_TYPE."); + return 0; } - return size; - } else if (type == CL_MEM_OBJECT_IMAGE2D) { - error = clGetImageInfo(mem, CL_IMAGE_WIDTH, sizeof(width), &width, NULL); - if (error) { - print_error(error, "clGetMemObjectInfo failed for CL_IMAGE_WIDTH."); - return 0; + + if (type == CL_MEM_OBJECT_BUFFER) + { + error = clGetMemObjectInfo(mem, CL_MEM_SIZE, sizeof(size), &size, NULL); + if (error) + { + print_error(error, "clGetMemObjectInfo failed for CL_MEM_SIZE."); + return 0; + } + return size; } - error = clGetImageInfo(mem, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL); - if (error) { - print_error(error, "clGetMemObjectInfo failed for CL_IMAGE_HEIGHT."); - return 0; + else if (type == CL_MEM_OBJECT_IMAGE2D) + { + error = + clGetImageInfo(mem, CL_IMAGE_WIDTH, sizeof(width), &width, NULL); + if (error) + { + print_error(error, "clGetMemObjectInfo failed for CL_IMAGE_WIDTH."); + return 0; + } + error = + clGetImageInfo(mem, CL_IMAGE_HEIGHT, sizeof(height), &height, NULL); + if (error) + { + print_error(error, + "clGetMemObjectInfo failed for CL_IMAGE_HEIGHT."); + return 0; + } + return width * height * 4 * sizeof(cl_uint); } - return width*height*4*sizeof(cl_uint); - } - log_error("Invalid CL_MEM_TYPE: %d\n", type); - return 0; + log_error("Invalid CL_MEM_TYPE: %d\n", type); + return 0; } - - diff --git a/test_conformance/allocations/allocation_utils.h b/test_conformance/allocations/allocation_utils.h index 2d165c1e05..241b139ced 100644 --- a/test_conformance/allocations/allocation_utils.h +++ b/test_conformance/allocations/allocation_utils.h @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -20,7 +20,9 @@ extern cl_uint checksum; -int check_allocation_error(cl_context context, cl_device_id device_id, int error, cl_command_queue *queue, cl_event *event = 0); +int check_allocation_error(cl_context context, cl_device_id device_id, + int error, cl_command_queue *queue, + cl_event *event = 0); double toMB(cl_ulong size_in); size_t get_actual_allocation_size(cl_mem mem); diff --git a/test_conformance/allocations/main.cpp b/test_conformance/allocations/main.cpp index 6ef83c680d..65d2699f3b 100644 --- a/test_conformance/allocations/main.cpp +++ b/test_conformance/allocations/main.cpp @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -31,66 +31,86 @@ int g_multiple_allocations = 0; int g_execute_kernel = 1; static size_t g_max_size; -static RandomSeed g_seed( gRandomSeed ); +static RandomSeed g_seed(gRandomSeed); cl_long g_max_individual_allocation_size; cl_long g_global_mem_size; cl_uint checksum; -static void printUsage( const char *execName ); +static void printUsage(const char *execName); -test_status init_cl( cl_device_id device ) { +test_status init_cl(cl_device_id device) +{ int error; - error = clGetDeviceInfo( device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(g_max_individual_allocation_size), &g_max_individual_allocation_size, NULL ); - if ( error ) { - print_error( error, "clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE"); + error = clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, + sizeof(g_max_individual_allocation_size), + &g_max_individual_allocation_size, NULL); + if (error) + { + print_error(error, + "clGetDeviceInfo failed for CL_DEVICE_MAX_MEM_ALLOC_SIZE"); return TEST_FAIL; } - error = clGetDeviceInfo( device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(g_global_mem_size), &g_global_mem_size, NULL ); - if ( error ) { - print_error( error, "clGetDeviceInfo failed for CL_DEVICE_GLOBAL_MEM_SIZE"); + error = + clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, + sizeof(g_global_mem_size), &g_global_mem_size, NULL); + if (error) + { + print_error(error, + "clGetDeviceInfo failed for CL_DEVICE_GLOBAL_MEM_SIZE"); return TEST_FAIL; } - log_info("Device reports CL_DEVICE_MAX_MEM_ALLOC_SIZE=%llu bytes (%gMB), CL_DEVICE_GLOBAL_MEM_SIZE=%llu bytes (%gMB).\n", - llu( g_max_individual_allocation_size ), toMB( g_max_individual_allocation_size ), - llu( g_global_mem_size ), toMB( g_global_mem_size ) ); + log_info("Device reports CL_DEVICE_MAX_MEM_ALLOC_SIZE=%llu bytes (%gMB), " + "CL_DEVICE_GLOBAL_MEM_SIZE=%llu bytes (%gMB).\n", + llu(g_max_individual_allocation_size), + toMB(g_max_individual_allocation_size), llu(g_global_mem_size), + toMB(g_global_mem_size)); - if( g_global_mem_size > (cl_ulong)SIZE_MAX ) + if (g_global_mem_size > (cl_ulong)SIZE_MAX) { g_global_mem_size = (cl_ulong)SIZE_MAX; } - if( g_max_individual_allocation_size > g_global_mem_size ) + if (g_max_individual_allocation_size > g_global_mem_size) { - log_error( "FAILURE: CL_DEVICE_MAX_MEM_ALLOC_SIZE (%llu) is greater than the CL_DEVICE_GLOBAL_MEM_SIZE (%llu)\n", - llu( g_max_individual_allocation_size ), llu( g_global_mem_size ) ); + log_error("FAILURE: CL_DEVICE_MAX_MEM_ALLOC_SIZE (%llu) is greater " + "than the CL_DEVICE_GLOBAL_MEM_SIZE (%llu)\n", + llu(g_max_individual_allocation_size), + llu(g_global_mem_size)); return TEST_FAIL; } - // We may need to back off the global_mem_size on unified memory devices to leave room for application and operating system code - // and associated data in the working set, so we dont start pathologically paging. - // Check to see if we are a unified memory device + // We may need to back off the global_mem_size on unified memory devices to + // leave room for application and operating system code and associated data + // in the working set, so we dont start pathologically paging. Check to see + // if we are a unified memory device cl_bool hasUnifiedMemory = CL_FALSE; - if( ( error = clGetDeviceInfo( device, CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof( hasUnifiedMemory ), &hasUnifiedMemory, NULL ) ) ) + if ((error = clGetDeviceInfo(device, CL_DEVICE_HOST_UNIFIED_MEMORY, + sizeof(hasUnifiedMemory), &hasUnifiedMemory, + NULL))) { - print_error( error, "clGetDeviceInfo failed for CL_DEVICE_HOST_UNIFIED_MEMORY"); + print_error(error, + "clGetDeviceInfo failed for CL_DEVICE_HOST_UNIFIED_MEMORY"); return TEST_FAIL; } // we share unified memory so back off to 1/2 the global memory size. - if( CL_TRUE == hasUnifiedMemory ) + if (CL_TRUE == hasUnifiedMemory) { - g_global_mem_size -= g_global_mem_size /2; - log_info( "Device shares memory with the host, so backing off the maximum combined allocation size to be %gMB to avoid rampant paging.\n", - toMB( g_global_mem_size ) ); + g_global_mem_size -= g_global_mem_size / 2; + log_info( + "Device shares memory with the host, so backing off the maximum " + "combined allocation size to be %gMB to avoid rampant paging.\n", + toMB(g_global_mem_size)); } else { - // Lets just use 60% of total available memory as framework/driver may not allow using all of it - // e.g. vram on GPU is used by window server and even for this test, we need some space for context, - // queue, kernel code on GPU. + // Lets just use 60% of total available memory as framework/driver may + // not allow using all of it e.g. vram on GPU is used by window server + // and even for this test, we need some space for context, queue, kernel + // code on GPU. g_global_mem_size *= 0.60; } /* Cap the allocation size as the global size was deduced */ @@ -99,15 +119,16 @@ test_status init_cl( cl_device_id device ) { g_max_individual_allocation_size = g_global_mem_size; } - if( gReSeed ) + if (gReSeed) { - g_seed = RandomSeed( gRandomSeed ); + g_seed = RandomSeed(gRandomSeed); } return TEST_PASS; } -int doTest( cl_device_id device, cl_context context, cl_command_queue queue, AllocType alloc_type ) +int doTest(cl_device_id device, cl_context context, cl_command_queue queue, + AllocType alloc_type) { int error; int failure_counts = 0; @@ -116,117 +137,141 @@ int doTest( cl_device_id device, cl_context context, cl_command_queue queue, All cl_mem mems[MAX_NUMBER_TO_ALLOCATE]; int number_of_mems_used; cl_ulong max_individual_allocation_size = g_max_individual_allocation_size; - cl_ulong global_mem_size = g_global_mem_size ; + cl_ulong global_mem_size = g_global_mem_size; + unsigned int number_of_work_itmes = 8192 * 32; const bool allocate_image = (alloc_type != BUFFER) && (alloc_type != BUFFER_NON_BLOCKING); - static const char* alloc_description[] = { - "buffer(s)", - "read-only image(s)", - "write-only image(s)", - "buffer(s)", - "read-only image(s)", - "write-only image(s)", + static const char *alloc_description[] = { + "buffer(s)", "read-only image(s)", "write-only image(s)", + "buffer(s)", "read-only image(s)", "write-only image(s)", }; // Skip image tests if we don't support images on the device if (allocate_image && checkForImageSupport(device)) { - log_info( "Can not test image allocation because device does not support images.\n" ); + log_info("Can not test image allocation because device does not " + "support images.\n"); return 0; } // This section was added in order to fix a bug in the test - // If CL_DEVICE_MAX_MEM_ALLOC_SIZE is much grater than CL_DEVICE_IMAGE2D_MAX_WIDTH * CL_DEVICE_IMAGE2D_MAX_HEIGHT - // The test will fail in image allocations as the size requested for the allocation will be much grater than the maximum size allowed for image + // If CL_DEVICE_MAX_MEM_ALLOC_SIZE is much grater than + // CL_DEVICE_IMAGE2D_MAX_WIDTH * CL_DEVICE_IMAGE2D_MAX_HEIGHT The test will + // fail in image allocations as the size requested for the allocation will + // be much grater than the maximum size allowed for image if (allocate_image) { size_t max_width, max_height; - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof( max_width ), &max_width, NULL ); - test_error_abort( error, "clGetDeviceInfo failed for CL_DEVICE_IMAGE2D_MAX_WIDTH" ); + error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, + sizeof(max_width), &max_width, NULL); + test_error_abort( + error, "clGetDeviceInfo failed for CL_DEVICE_IMAGE2D_MAX_WIDTH"); - error = clGetDeviceInfo( device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof( max_height ), &max_height, NULL ); - test_error_abort( error, "clGetDeviceInfo failed for CL_DEVICE_IMAGE2D_MAX_HEIGHT" ); + error = clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, + sizeof(max_height), &max_height, NULL); + test_error_abort( + error, "clGetDeviceInfo failed for CL_DEVICE_IMAGE2D_MAX_HEIGHT"); - cl_ulong max_image2d_size = (cl_ulong)max_height * max_width * 4 * sizeof(cl_uint); + cl_ulong max_image2d_size = + (cl_ulong)max_height * max_width * 4 * sizeof(cl_uint); - if( max_individual_allocation_size > max_image2d_size ) + if (max_individual_allocation_size > max_image2d_size) { max_individual_allocation_size = max_image2d_size; } } - // Pick the baseline size based on whether we are doing a single large or multiple allocations - g_max_size = g_multiple_allocations ? (size_t)global_mem_size : (size_t)max_individual_allocation_size; + // Pick the baseline size based on whether we are doing a single large or + // multiple allocations + g_max_size = g_multiple_allocations + ? (size_t)global_mem_size + : (size_t)max_individual_allocation_size; // Adjust based on the percentage - if( g_reduction_percentage != 100 ) + if (g_reduction_percentage != 100) { - log_info( "NOTE: reducing max allocations to %d%%.\n", g_reduction_percentage ); - g_max_size = (size_t)( (double)g_max_size * (double)g_reduction_percentage / 100.0 ); + log_info("NOTE: reducing max allocations to %d%%.\n", + g_reduction_percentage); + g_max_size = (size_t)((double)g_max_size + * (double)g_reduction_percentage / 100.0); + number_of_work_itmes = 8192 * 2; } // Round to nearest MB. g_max_size &= (size_t)(0xFFFFFFFFFF00000ULL); - log_info( "** Target allocation size (rounded to nearest MB) is: %llu bytes (%gMB).\n", llu( g_max_size ), toMB( g_max_size ) ); - log_info( "** Allocating %s to size %gMB.\n", alloc_description[alloc_type], toMB( g_max_size ) ); + log_info("** Target allocation size (rounded to nearest MB) is: %llu bytes " + "(%gMB).\n", + llu(g_max_size), toMB(g_max_size)); + log_info("** Allocating %s to size %gMB.\n", alloc_description[alloc_type], + toMB(g_max_size)); - for( int count = 0; count < g_repetition_count; count++ ) + for (int count = 0; count < g_repetition_count; count++) { current_test_size = g_max_size; error = FAILED_TOO_BIG; - log_info( " => Allocation %d\n", count + 1 ); + log_info(" => Allocation %d\n", count + 1); - while( ( error == FAILED_TOO_BIG ) && ( current_test_size > g_max_size / 8 ) ) + while ((error == FAILED_TOO_BIG) + && (current_test_size > g_max_size / 8)) { // Reset our checksum for each allocation checksum = 0; // Do the allocation - error = allocate_size( context, &queue, device, g_multiple_allocations, current_test_size, alloc_type, - mems, &number_of_mems_used, &final_size, g_write_allocations, g_seed ); + error = allocate_size(context, &queue, device, + g_multiple_allocations, current_test_size, + alloc_type, mems, &number_of_mems_used, + &final_size, g_write_allocations, g_seed); // If we succeeded and we're supposed to execute a kernel, do so. - if( error == SUCCEEDED && g_execute_kernel ) + if (error == SUCCEEDED && g_execute_kernel) { - log_info( "\tExecuting kernel with memory objects.\n" ); - error = execute_kernel( context, &queue, device, alloc_type, mems, number_of_mems_used, - g_write_allocations ); + log_info("\tExecuting kernel with memory objects.\n"); + error = + execute_kernel(context, &queue, device, alloc_type, mems, + number_of_mems_used, g_write_allocations, + number_of_work_itmes); } - // If we failed to allocate more than 1/8th of the requested amount return a failure. - if( final_size < (size_t)g_max_size / 8 ) + // If we failed to allocate more than 1/8th of the requested amount + // return a failure. + if (final_size < (size_t)g_max_size / 8) { - log_error( "===> Allocation %d failed to allocate more than 1/8th of the requested size.\n", count + 1 ); + log_error("===> Allocation %d failed to allocate more than " + "1/8th of the requested size.\n", + count + 1); failure_counts++; } // Clean up. - for( int i = 0; i < number_of_mems_used; i++ ) + for (int i = 0; i < number_of_mems_used; i++) { - clReleaseMemObject( mems[i] ); + clReleaseMemObject(mems[i]); } - if( error == FAILED_ABORT ) + if (error == FAILED_ABORT) { - log_error( " => Allocation %d failed.\n", count + 1 ); + log_error(" => Allocation %d failed.\n", count + 1); failure_counts++; } - if( error == FAILED_TOO_BIG ) + if (error == FAILED_TOO_BIG) { current_test_size -= g_max_size / 16; - log_info( "\tFailed at this size; trying a smaller size of %gMB.\n", toMB( current_test_size ) ); + log_info( + "\tFailed at this size; trying a smaller size of %gMB.\n", + toMB(current_test_size)); } } - if( error == SUCCEEDED && current_test_size == g_max_size ) + if (error == SUCCEEDED && current_test_size == g_max_size) { log_info("\tPASS: Allocation succeeded.\n"); } - else if( error == SUCCEEDED && current_test_size > g_max_size / 8 ) + else if (error == SUCCEEDED && current_test_size > g_max_size / 8) { log_info("\tPASS: Allocation succeeded at reduced size.\n"); } @@ -240,41 +285,47 @@ int doTest( cl_device_id device, cl_context context, cl_command_queue queue, All return failure_counts; } -int test_buffer(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_buffer(cl_device_id device, cl_context context, cl_command_queue queue, + int num_elements) { - return doTest( device, context, queue, BUFFER ); + return doTest(device, context, queue, BUFFER); } -int test_image2d_read(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_image2d_read(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest( device, context, queue, IMAGE_READ ); + return doTest(device, context, queue, IMAGE_READ); } -int test_image2d_write(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_image2d_write(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest( device, context, queue, IMAGE_WRITE ); + return doTest(device, context, queue, IMAGE_WRITE); } -int test_buffer_non_blocking(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_buffer_non_blocking(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest( device, context, queue, BUFFER_NON_BLOCKING ); + return doTest(device, context, queue, BUFFER_NON_BLOCKING); } -int test_image2d_read_non_blocking(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_image2d_read_non_blocking(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest( device, context, queue, IMAGE_READ_NON_BLOCKING ); + return doTest(device, context, queue, IMAGE_READ_NON_BLOCKING); } -int test_image2d_write_non_blocking(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) +int test_image2d_write_non_blocking(cl_device_id device, cl_context context, + cl_command_queue queue, int num_elements) { - return doTest( device, context, queue, IMAGE_WRITE_NON_BLOCKING ); + return doTest(device, context, queue, IMAGE_WRITE_NON_BLOCKING); } test_definition test_list[] = { - ADD_TEST( buffer ), - ADD_TEST( image2d_read ), - ADD_TEST( image2d_write ), - ADD_TEST( buffer_non_blocking ), - ADD_TEST( image2d_read_non_blocking ), - ADD_TEST( image2d_write_non_blocking ), + ADD_TEST(buffer), + ADD_TEST(image2d_read), + ADD_TEST(image2d_write), + ADD_TEST(buffer_non_blocking), + ADD_TEST(image2d_read_non_blocking), + ADD_TEST(image2d_write_non_blocking), }; -const int test_num = ARRAY_SIZE( test_list ); +const int test_num = ARRAY_SIZE(test_list); int main(int argc, const char *argv[]) { @@ -287,11 +338,11 @@ int main(int argc, const char *argv[]) return 1; } - const char ** argList = (const char **)calloc( argc, sizeof( char*) ); + const char **argList = (const char **)calloc(argc, sizeof(char *)); - if( NULL == argList ) + if (NULL == argList) { - log_error( "Failed to allocate memory for argList array.\n" ); + log_error("Failed to allocate memory for argList array.\n"); return 1; } @@ -299,38 +350,40 @@ int main(int argc, const char *argv[]) size_t argCount = 1; // Parse arguments - for( int i = 1; i < argc; i++ ) + for (int i = 1; i < argc; i++) { - if( strcmp( argv[i], "multiple" ) == 0 ) + if (strcmp(argv[i], "multiple") == 0) g_multiple_allocations = 1; - else if( strcmp( argv[i], "single" ) == 0 ) + else if (strcmp(argv[i], "single") == 0) g_multiple_allocations = 0; - else if( ( r = (int)strtol( argv[i], &endPtr, 10 ) ) && ( endPtr != argv[i] ) && ( *endPtr == 0 ) ) + else if ((r = (int)strtol(argv[i], &endPtr, 10)) && (endPtr != argv[i]) + && (*endPtr == 0)) { - // By spec, that means the entire string was an integer, so take it as a repetition count + // By spec, that means the entire string was an integer, so take it + // as a repetition count g_repetition_count = r; } - else if( strchr( argv[i], '%' ) != NULL ) + else if (strchr(argv[i], '%') != NULL) { // Reduction percentage (let strtol ignore the percentage) - g_reduction_percentage = (int)strtol( argv[i], NULL, 10 ); + g_reduction_percentage = (int)strtol(argv[i], NULL, 10); } - else if( strcmp( argv[i], "do_not_force_fill" ) == 0 ) + else if (strcmp(argv[i], "do_not_force_fill") == 0) { g_write_allocations = 0; } - else if( strcmp( argv[i], "do_not_execute" ) == 0 ) + else if (strcmp(argv[i], "do_not_execute") == 0) { g_execute_kernel = 0; } - else if ( strcmp( argv[i], "--help" ) == 0 || strcmp( argv[i], "-h" ) == 0 ) + else if (strcmp(argv[i], "--help") == 0 || strcmp(argv[i], "-h") == 0) { - printUsage( argv[0] ); + printUsage(argv[0]); free(argList); return -1; } @@ -342,35 +395,42 @@ int main(int argc, const char *argv[]) } } - int ret = runTestHarnessWithCheck( argCount, argList, test_num, test_list, false, 0, init_cl ); + int ret = runTestHarnessWithCheck(argCount, argList, test_num, test_list, + false, 0, init_cl); free(argList); return ret; } -void printUsage( const char *execName ) +void printUsage(const char *execName) { - const char *p = strrchr( execName, '/' ); - if( p != NULL ) - execName = p + 1; - - log_info( "Usage: %s [options] [test_names]\n", execName ); - log_info( "Options:\n" ); - log_info( "\trandomize - Uses random seed\n" ); - log_info( "\tsingle - Tests using a single allocation as large as possible\n" ); - log_info( "\tmultiple - Tests using as many allocations as possible\n" ); - log_info( "\n" ); - log_info( "\tnumReps - Optional integer specifying the number of repetitions to run and average the result (defaults to 1)\n" ); - log_info( "\treduction%% - Optional integer, followed by a %% sign, that acts as a multiplier for the target amount of memory.\n" ); - log_info( "\t Example: target amount of 512MB and a reduction of 75%% will result in a target of 384MB.\n" ); - log_info( "\n" ); - log_info( "\tdo_not_force_fill - Disable explicitly write data to all memory objects after creating them.\n" ); - log_info( "\t Without this, the kernel execution can not verify its checksum.\n" ); - log_info( "\tdo_not_execute - Disable executing a kernel that accesses all of the memory objects.\n" ); - log_info( "\n" ); - log_info( "Test names (Allocation Types):\n" ); - for( int i = 0; i < test_num; i++ ) + const char *p = strrchr(execName, '/'); + if (p != NULL) execName = p + 1; + + log_info("Usage: %s [options] [test_names]\n", execName); + log_info("Options:\n"); + log_info("\trandomize - Uses random seed\n"); + log_info( + "\tsingle - Tests using a single allocation as large as possible\n"); + log_info("\tmultiple - Tests using as many allocations as possible\n"); + log_info("\n"); + log_info("\tnumReps - Optional integer specifying the number of " + "repetitions to run and average the result (defaults to 1)\n"); + log_info("\treduction%% - Optional integer, followed by a %% sign, that " + "acts as a multiplier for the target amount of memory.\n"); + log_info("\t Example: target amount of 512MB and a reduction " + "of 75%% will result in a target of 384MB.\n"); + log_info("\n"); + log_info("\tdo_not_force_fill - Disable explicitly write data to all " + "memory objects after creating them.\n"); + log_info("\t Without this, the kernel execution can not " + "verify its checksum.\n"); + log_info("\tdo_not_execute - Disable executing a kernel that accesses all " + "of the memory objects.\n"); + log_info("\n"); + log_info("Test names (Allocation Types):\n"); + for (int i = 0; i < test_num; i++) { - log_info( "\t%s\n", test_list[i].name ); + log_info("\t%s\n", test_list[i].name); } } diff --git a/test_conformance/allocations/testBase.h b/test_conformance/allocations/testBase.h index b48efe5188..d320907522 100644 --- a/test_conformance/allocations/testBase.h +++ b/test_conformance/allocations/testBase.h @@ -1,6 +1,6 @@ // // Copyright (c) 2017 The Khronos Group Inc. -// +// // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at @@ -39,9 +39,10 @@ #define FAILED_CORRUPTED_QUEUE -2 #define FAILED_ABORT -1 #define FAILED_TOO_BIG 1 -// On Windows macro `SUCCEEDED' is defined in `WinError.h'. It causes compiler warnings. Let us avoid them. -#if defined( _WIN32 ) && defined( SUCCEEDED ) - #undef SUCCEEDED +// On Windows macro `SUCCEEDED' is defined in `WinError.h'. It causes compiler +// warnings. Let us avoid them. +#if defined(_WIN32) && defined(SUCCEEDED) +#undef SUCCEEDED #endif #define SUCCEEDED 0 @@ -55,11 +56,16 @@ enum AllocType IMAGE_WRITE_NON_BLOCKING, }; -#define test_error_abort(errCode,msg) test_error_ret_abort(errCode,msg,errCode) -#define test_error_ret_abort(errCode,msg,retValue) { if( errCode != CL_SUCCESS ) { print_error( errCode, msg ); return FAILED_ABORT ; } } +#define test_error_abort(errCode, msg) \ + test_error_ret_abort(errCode, msg, errCode) +#define test_error_ret_abort(errCode, msg, retValue) \ + { \ + if (errCode != CL_SUCCESS) \ + { \ + print_error(errCode, msg); \ + return FAILED_ABORT; \ + } \ + } #endif // _testBase_h - - -