From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 9627 invoked by alias); 10 Dec 2015 12:06:02 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 9549 invoked by uid 89); 10 Dec 2015 12:06:02 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=0.6 required=5.0 tests=AWL,BAYES_50,KAM_STOCKGEN,RCVD_IN_DNSWL_LOW,SPF_PASS autolearn=no version=3.3.2 X-HELO: mx2.suse.de Received: from mx2.suse.de (HELO mx2.suse.de) (195.135.220.15) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with (CAMELLIA256-SHA encrypted) ESMTPS; Thu, 10 Dec 2015 12:05:58 +0000 Received: from relay2.suse.de (charybdis-ext.suse.de [195.135.220.254]) by mx2.suse.de (Postfix) with ESMTP id 7F03EAC3B for ; Thu, 10 Dec 2015 12:05:55 +0000 (UTC) Subject: Re: [hsa 3/10] HSA libgomp plugin [part 1/2] To: gcc-patches@gcc.gnu.org References: <20151207111758.GA24234@virgil.suse.cz> <20151207112049.GD24234@virgil.suse.cz> <20151209121555.GQ5675@tucnak.redhat.com> From: =?UTF-8?Q?Martin_Li=c5=a1ka?= Message-ID: <56696AA2.4070305@suse.cz> Date: Thu, 10 Dec 2015 12:06:00 -0000 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:38.0) Gecko/20100101 Thunderbird/38.3.0 MIME-Version: 1.0 In-Reply-To: <20151209121555.GQ5675@tucnak.redhat.com> Content-Type: multipart/mixed; boundary="------------000200030807060004070004" X-IsSubscribed: yes X-SW-Source: 2015-12/txt/msg01114.txt.bz2 This is a multi-part message in MIME format. --------------000200030807060004070004 Content-Type: text/plain; charset=windows-1252 Content-Transfer-Encoding: 7bit Content-length: 5162 On 12/09/2015 01:15 PM, Jakub Jelinek wrote: > On Mon, Dec 07, 2015 at 12:20:49PM +0100, Martin Jambor wrote: >> +#include >> +#include >> +#include >> +#include >> +#include "libgomp-plugin.h" >> +#include "gomp-constants.h" >> +#include "hsa.h" >> +#include "hsa_ext_finalize.h" > > If these 2 headers are coming from outside of gcc, better use <> for them > instead of "", and better place them above the libgomp specific ones. > >> +#include "dlfcn.h" > > Ditto. > >> +/* Flag to decide whether print to stderr information about what is going on. >> + Set in init_debug depending on environment variables. */ >> + >> +static bool debug; >> + >> +/* Flag to decide if the runtime should suppress a possible fallback to host >> + execution. */ >> + >> +static bool suppress_host_fallback; >> + >> +/* Initialize debug and suppress_host_fallback according to the environment. */ >> + >> +static void >> +init_enviroment_variables (void) >> +{ >> + if (getenv ("HSA_DEBUG")) >> + debug = true; >> + else >> + debug = false; >> + >> + if (getenv ("HSA_SUPPRESS_HOST_FALLBACK")) >> + suppress_host_fallback = true; >> + else >> + suppress_host_fallback = false; > > Wonder whether we want these custom env vars in suid programs, allowing > users to change behavior might be undesirable. So perhaps use secure_getenv > instead of getenv if found by configure? >> + >> + const char* hsa_error; > > Space before * instead of after it (multiple spots). > >> + hsa_status_string (status, &hsa_error); >> + >> + unsigned l = strlen (hsa_error); >> + >> + char *err = GOMP_PLUGIN_malloc (sizeof (char) * l); > > sizeof (char) is 1, please don't multiply by it, that is just confusing. > It makes sense in macros where you don't know what the type really is, but > not here. > >> + memcpy (err, hsa_error, l - 1); >> + err[l] = '\0'; >> + >> + fprintf (stderr, "HSA warning: %s (%s)\n", str, err); > > Why do you allocate err at all, if you free it right away? Can't you use > hsa_error directly instead? >> + >> + free (err); > >> +static void >> +hsa_fatal (const char *str, hsa_status_t status) >> +{ >> + const char* hsa_error; >> + hsa_status_string (status, &hsa_error); >> + GOMP_PLUGIN_fatal ("HSA fatal error: %s (%s)", str, hsa_error); > > Like you do e.g. here? > >> +/* Callback of dispatch queues to report errors. */ >> + >> +static void >> +queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)), > > Missing space before (. >> +/* Callback of hsa_agent_iterate_regions. Determine if a memory REGION can be >> + used for kernarg allocations and if so write it to the memory pointed to by >> + DATA and break the query. */ >> + >> +static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data) > > Missing newline between return type and function name. > >> + hsa_region_t* ret = (hsa_region_t*) data; > > Two spots with wrong formatting above. >> +{ >> + if (agent->first_module) >> + agent->first_module->prev = module; > > Wrong indentation. > >> + unsigned size = end - start; >> + char *buf = (char *) malloc (size); >> + memcpy (buf, start, size); > > malloc could return NULL and you crash. Should this use GOMP_PLUGIN_malloc > instead? > >> + struct GOMP_hsa_kernel_dispatch *shadow = GOMP_PLUGIN_malloc_cleared >> + (sizeof (struct GOMP_hsa_kernel_dispatch)); > > Formatting, = should go on the next line, and if even that doesn't help, > maybe use sizeof (*shadow)? >> + >> + shadow->queue = agent->command_q; >> + shadow->omp_data_memory = omp_data_size > 0 >> + ? GOMP_PLUGIN_malloc (omp_data_size) : NULL; > > = should go on the next line. > >> + unsigned dispatch_count = kernel->dependencies_count; >> + shadow->kernel_dispatch_count = dispatch_count; >> + >> + shadow->children_dispatches = GOMP_PLUGIN_malloc >> + (dispatch_count * sizeof (struct GOMP_hsa_kernel_dispatch *)); > > Likewise. >> +indent_stream (FILE *f, unsigned indent) >> +{ >> + for (int i = 0; i < indent; i++) >> + fputc (' ', f); > > Perhaps fprintf (f, "%*s", indent, ""); > instead? > >> + struct GOMP_hsa_kernel_dispatch *shadow = create_single_kernel_dispatch >> + (kernel, omp_data_size); > > Formatting issues again. > >> + shadow->omp_num_threads = 64; >> + shadow->debug = 0; >> + shadow->omp_level = kernel->gridified_kernel_p ? 1 : 0; >> + >> + /* Create kernel dispatch data structures. We do not allow to have >> + a kernel dispatch with depth bigger than one. */ >> + for (unsigned i = 0; i < kernel->dependencies_count; i++) >> + { >> + struct kernel_info *dependency = get_kernel_for_agent >> + (kernel->agent, kernel->dependencies[i]); >> + shadow->children_dispatches[i] = create_single_kernel_dispatch >> + (dependency, omp_data_size); >> + shadow->children_dispatches[i]->queue = > > Never put = at the end of line. >> + kernel->agent->kernel_dispatch_command_q; > > Jakub > Hello. There's a new version of the patch that changes many GNU coding style issues, trimming of last character in hsa_warn is removed. Martin --------------000200030807060004070004 Content-Type: text/x-patch; name="0001-Clean-up-plugin-hsa.c-file.patch" Content-Transfer-Encoding: 7bit Content-Disposition: attachment; filename="0001-Clean-up-plugin-hsa.c-file.patch" Content-length: 15290 >From a289f04ce17d8f3a0c30ce87b029854736f82a16 Mon Sep 17 00:00:00 2001 From: marxin Date: Wed, 9 Dec 2015 14:06:55 +0100 Subject: [PATCH 1/2] Clean-up plugin-hsa.c file libgomp/ChangeLog: 2015-12-09 Martin Liska * plugin/plugin-hsa.c (hsa_warn): Change output string format. (hsa_fatal): Dtto. (struct agent_info): Fix GNU coding style. (suitable_hsa_agent_p): Likewise. (get_kernarg_memory_region): Likewise. (add_module_to_agent): Likewise. (add_shared_library): Use GOMP_PLUGIN_malloc instead of malloc. (create_and_finalize_hsa_program): Fix GNU coding style. (create_single_kernel_dispatch): Dtto. (release_kernel_dispatch): Dtto. (init_single_kernel): Dtto. (indent_stream): Use printf. (print_kernel_dispatch): Fix GNU coding style. (create_kernel_dispatch): Dtto. (GOMP_OFFLOAD_run): Dtto. --- libgomp/plugin/plugin-hsa.c | 144 ++++++++++++++++++++++---------------------- 1 file changed, 71 insertions(+), 73 deletions(-) diff --git a/libgomp/plugin/plugin-hsa.c b/libgomp/plugin/plugin-hsa.c index b132954..46b9c03 100644 --- a/libgomp/plugin/plugin-hsa.c +++ b/libgomp/plugin/plugin-hsa.c @@ -31,11 +31,11 @@ #include #include #include +#include +#include +#include #include "libgomp-plugin.h" #include "gomp-constants.h" -#include "hsa.h" -#include "hsa_ext_finalize.h" -#include "dlfcn.h" /* Part of the libgomp plugin interface. Return the name of the accelerator, which is "hsa". */ @@ -128,18 +128,10 @@ hsa_warn (const char *str, hsa_status_t status) if (!debug) return; - const char* hsa_error; + const char *hsa_error; hsa_status_string (status, &hsa_error); - unsigned l = strlen (hsa_error); - - char *err = GOMP_PLUGIN_malloc (sizeof (char) * l); - memcpy (err, hsa_error, l - 1); - err[l] = '\0'; - - fprintf (stderr, "HSA warning: %s (%s)\n", str, err); - - free (err); + fprintf (stderr, "HSA warning: %s\nRuntime message: %s", str, hsa_error); } /* Report a fatal error STR together with the HSA error corresponding to STATUS @@ -148,9 +140,10 @@ hsa_warn (const char *str, hsa_status_t status) static void hsa_fatal (const char *str, hsa_status_t status) { - const char* hsa_error; + const char *hsa_error; hsa_status_string (status, &hsa_error); - GOMP_PLUGIN_fatal ("HSA fatal error: %s (%s)", str, hsa_error); + GOMP_PLUGIN_fatal ("HSA fatal error: %s\nRuntime message: %s", str, + hsa_error); } struct hsa_kernel_description @@ -258,9 +251,9 @@ struct agent_info /* The HSA ISA of this agent. */ hsa_isa_t isa; /* Command queue of the agent. */ - hsa_queue_t* command_q; + hsa_queue_t *command_q; /* Kernel from kernel dispatch command queue. */ - hsa_queue_t* kernel_dispatch_command_q; + hsa_queue_t *kernel_dispatch_command_q; /* The HSA memory region from which to allocate kernel arguments. */ hsa_region_t kernarg_region; @@ -332,8 +325,8 @@ static bool suitable_hsa_agent_p (hsa_agent_t agent) { hsa_device_type_t device_type; - hsa_status_t status = hsa_agent_get_info (agent, HSA_AGENT_INFO_DEVICE, - &device_type); + hsa_status_t status + = hsa_agent_get_info (agent, HSA_AGENT_INFO_DEVICE, &device_type); if (status != HSA_STATUS_SUCCESS || device_type != HSA_DEVICE_TYPE_GPU) return false; @@ -410,8 +403,9 @@ init_hsa_context (void) /* Callback of dispatch queues to report errors. */ static void -queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)), - void* data __attribute__ ((unused))) +queue_callback (hsa_status_t status, + hsa_queue_t *queue __attribute__ ((unused)), + void *data __attribute__ ((unused))) { hsa_fatal ("Asynchronous queue error", status); } @@ -420,7 +414,8 @@ queue_callback(hsa_status_t status, hsa_queue_t* queue __attribute__ ((unused)), used for kernarg allocations and if so write it to the memory pointed to by DATA and break the query. */ -static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data) +static hsa_status_t +get_kernarg_memory_region (hsa_region_t region, void *data) { hsa_status_t status; hsa_region_segment_t segment; @@ -437,7 +432,7 @@ static hsa_status_t get_kernarg_memory_region (hsa_region_t region, void* data) return status; if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) { - hsa_region_t* ret = (hsa_region_t*) data; + hsa_region_t *ret = (hsa_region_t *) data; *ret = region; return HSA_STATUS_INFO_BREAK; } @@ -528,7 +523,7 @@ static void add_module_to_agent (struct agent_info *agent, struct module_info *module) { if (agent->first_module) - agent->first_module->prev = module; + agent->first_module->prev = module; module->next = agent->first_module; module->prev = NULL; agent->first_module = module; @@ -653,12 +648,12 @@ add_shared_library (const char *file_name, struct agent_info *agent) return NULL; unsigned size = end - start; - char *buf = (char *) malloc (size); + char *buf = (char *) GOMP_PLUGIN_malloc (size); memcpy (buf, start, size); library = GOMP_PLUGIN_malloc (sizeof (struct agent_info)); library->file_name = (char *) GOMP_PLUGIN_malloc - ((strlen (file_name) + 1) * sizeof (char)); + ((strlen (file_name) + 1)); strcpy (library->file_name, file_name); library->image = (hsa_ext_module_t) buf; @@ -746,11 +741,11 @@ create_and_finalize_hsa_program (struct agent_info *agent) hsa_ext_control_directives_t control_directives; memset (&control_directives, 0, sizeof (control_directives)); hsa_code_object_t code_object; - status = hsa_ext_program_finalize(prog_handle, agent->isa, - HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO, - control_directives, "", - HSA_CODE_OBJECT_TYPE_PROGRAM, - &code_object); + status = hsa_ext_program_finalize (prog_handle, agent->isa, + HSA_EXT_FINALIZER_CALL_CONVENTION_AUTO, + control_directives, "", + HSA_CODE_OBJECT_TYPE_PROGRAM, + &code_object); if (status != HSA_STATUS_SUCCESS) { hsa_warn ("Finalization of the HSA program failed", status); @@ -760,8 +755,9 @@ create_and_finalize_hsa_program (struct agent_info *agent) HSA_DEBUG ("Finalization done\n"); hsa_ext_program_destroy (prog_handle); - status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, - "", &agent->executable); + status + = hsa_executable_create (HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, + "", &agent->executable); if (status != HSA_STATUS_SUCCESS) hsa_fatal ("Could not create HSA executable", status); @@ -773,8 +769,9 @@ create_and_finalize_hsa_program (struct agent_info *agent) { struct global_var_info *var; var = &module->image_desc->global_variables[i]; - status = hsa_executable_global_variable_define - (agent->executable, var->name, var->address); + status + = hsa_executable_global_variable_define (agent->executable, + var->name, var->address); HSA_DEBUG ("Defining global variable: %s, address: %p\n", var->name, var->address); @@ -787,11 +784,11 @@ create_and_finalize_hsa_program (struct agent_info *agent) module = module->next; } - status = hsa_executable_load_code_object(agent->executable, agent->id, - code_object, ""); + status = hsa_executable_load_code_object (agent->executable, agent->id, + code_object, ""); if (status != HSA_STATUS_SUCCESS) hsa_fatal ("Could not add a code object to the HSA executable", status); - status = hsa_executable_freeze(agent->executable, ""); + status = hsa_executable_freeze (agent->executable, ""); if (status != HSA_STATUS_SUCCESS) hsa_fatal ("Could not freeze the HSA executable", status); @@ -818,17 +815,17 @@ create_single_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size) { struct agent_info *agent = kernel->agent; - struct GOMP_hsa_kernel_dispatch *shadow = GOMP_PLUGIN_malloc_cleared - (sizeof (struct GOMP_hsa_kernel_dispatch)); + struct GOMP_hsa_kernel_dispatch *shadow + = GOMP_PLUGIN_malloc_cleared (sizeof (struct GOMP_hsa_kernel_dispatch)); shadow->queue = agent->command_q; - shadow->omp_data_memory = omp_data_size > 0 - ? GOMP_PLUGIN_malloc (omp_data_size) : NULL; + shadow->omp_data_memory + = omp_data_size > 0 ? GOMP_PLUGIN_malloc (omp_data_size) : NULL; unsigned dispatch_count = kernel->dependencies_count; shadow->kernel_dispatch_count = dispatch_count; - shadow->children_dispatches = GOMP_PLUGIN_malloc - (dispatch_count * sizeof (struct GOMP_hsa_kernel_dispatch *)); + shadow->children_dispatches + = GOMP_PLUGIN_malloc (dispatch_count * sizeof (shadow)); shadow->object = kernel->object; @@ -841,9 +838,9 @@ create_single_kernel_dispatch (struct kernel_info *kernel, shadow->private_segment_size = kernel->private_segment_size; shadow->group_segment_size = kernel->group_segment_size; - status = hsa_memory_allocate - (agent->kernarg_region, kernel->kernarg_segment_size, - &shadow->kernarg_address); + status + = hsa_memory_allocate (agent->kernarg_region, kernel->kernarg_segment_size, + &shadow->kernarg_address); if (status != HSA_STATUS_SUCCESS) hsa_fatal ("Could not allocate memory for HSA kernel arguments", status); @@ -855,8 +852,8 @@ create_single_kernel_dispatch (struct kernel_info *kernel, static void release_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *shadow) { - HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n", - shadow, shadow->debug, (void *)shadow->debug); + HSA_DEBUG ("Released kernel dispatch: %p has value: %lu (%p)\n", shadow, + shadow->debug, (void *) shadow->debug); hsa_memory_free (shadow->kernarg_address); @@ -890,8 +887,10 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size) goto failure; } HSA_DEBUG ("Located kernel %s\n", kernel->name); - status = hsa_executable_symbol_get_info - (kernel_symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &kernel->object); + status + = hsa_executable_symbol_get_info (kernel_symbol, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, + &kernel->object); if (status != HSA_STATUS_SUCCESS) hsa_fatal ("Could not extract a kernel object from its symbol", status); status = hsa_executable_symbol_get_info @@ -927,8 +926,8 @@ init_single_kernel (struct kernel_info *kernel, unsigned *max_omp_data_size) for (unsigned i = 0; i < kernel->dependencies_count; i++) { - struct kernel_info *dependency = get_kernel_for_agent - (agent, kernel->dependencies[i]); + struct kernel_info *dependency + = get_kernel_for_agent (agent, kernel->dependencies[i]); if (dependency == NULL) { @@ -959,8 +958,7 @@ failure: static void indent_stream (FILE *f, unsigned indent) { - for (int i = 0; i < indent; i++) - fputc (' ', f); + fprintf (f, "%*s", indent, ""); } /* Dump kernel DISPATCH data structure and indent it by INDENT spaces. */ @@ -995,7 +993,7 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden fprintf (stderr, "\n"); for (unsigned i = 0; i < dispatch->kernel_dispatch_count; i++) - print_kernel_dispatch (dispatch->children_dispatches[i], indent + 2); + print_kernel_dispatch (dispatch->children_dispatches[i], indent + 2); } /* Create kernel dispatch data structure for a KERNEL and all its @@ -1004,8 +1002,8 @@ print_kernel_dispatch (struct GOMP_hsa_kernel_dispatch *dispatch, unsigned inden static struct GOMP_hsa_kernel_dispatch * create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size) { - struct GOMP_hsa_kernel_dispatch *shadow = create_single_kernel_dispatch - (kernel, omp_data_size); + struct GOMP_hsa_kernel_dispatch *shadow + = create_single_kernel_dispatch (kernel, omp_data_size); shadow->omp_num_threads = 64; shadow->debug = 0; shadow->omp_level = kernel->gridified_kernel_p ? 1 : 0; @@ -1014,12 +1012,12 @@ create_kernel_dispatch (struct kernel_info *kernel, unsigned omp_data_size) a kernel dispatch with depth bigger than one. */ for (unsigned i = 0; i < kernel->dependencies_count; i++) { - struct kernel_info *dependency = get_kernel_for_agent - (kernel->agent, kernel->dependencies[i]); - shadow->children_dispatches[i] = create_single_kernel_dispatch - (dependency, omp_data_size); - shadow->children_dispatches[i]->queue = - kernel->agent->kernel_dispatch_command_q; + struct kernel_info *dependency + = get_kernel_for_agent (kernel->agent, kernel->dependencies[i]); + shadow->children_dispatches[i] + = create_single_kernel_dispatch (dependency, omp_data_size); + shadow->children_dispatches[i]->queue + = kernel->agent->kernel_dispatch_command_q; shadow->children_dispatches[i]->omp_level = 1; } @@ -1146,7 +1144,7 @@ failure: FN_PTR which must point to a kernel_info structure. */ void -GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args) +GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void **args) { struct kernel_info *kernel = (struct kernel_info *) fn_ptr; struct agent_info *agent = kernel->agent; @@ -1166,8 +1164,8 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args) if (!kernel->initialized) GOMP_PLUGIN_fatal ("Called kernel must be initialized"); - struct GOMP_hsa_kernel_dispatch *shadow = create_kernel_dispatch - (kernel, kernel->max_omp_data_size); + struct GOMP_hsa_kernel_dispatch *shadow + = create_kernel_dispatch (kernel, kernel->max_omp_data_size); if (debug) { @@ -1179,16 +1177,16 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args) HSA_DEBUG ("Got AQL index %llu\n", (long long int) index); /* Wait until the queue is not full before writing the packet. */ - while (index - hsa_queue_load_read_index_acquire(agent->command_q) + while (index - hsa_queue_load_read_index_acquire (agent->command_q) >= agent->command_q->size) ; hsa_kernel_dispatch_packet_t *packet; - packet = ((hsa_kernel_dispatch_packet_t*) agent->command_q->base_address) - + index % agent->command_q->size; + packet = ((hsa_kernel_dispatch_packet_t *) agent->command_q->base_address) + + index % agent->command_q->size; - memset (((uint8_t *)packet) + 4, 0, sizeof (*packet) - 4); - packet->setup |= (uint16_t) 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + memset (((uint8_t *) packet) + 4, 0, sizeof (*packet) - 4); + packet->setup |= (uint16_t) 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; packet->grid_size_x = kla->gdims[0]; uint32_t wgs = kla->wdims[0]; if (wgs == 0) @@ -1223,7 +1221,7 @@ GOMP_OFFLOAD_run (int n, void *fn_ptr, void *vars, void** args) HSA_DEBUG ("Going to dispatch kernel %s\n", kernel->name); - __atomic_store_n ((uint16_t*)(&packet->header), header, __ATOMIC_RELEASE); + __atomic_store_n ((uint16_t *) (&packet->header), header, __ATOMIC_RELEASE); hsa_signal_store_release (agent->command_q->doorbell_signal, index); /* TODO: GPU agents in Carrizo APUs cannot properly update L2 cache for -- 2.6.3 --------------000200030807060004070004--