Add Support for Injecting Buffers (#359)

* initial support for injecting buffers

* initial support for injecting images

* cleanup, updated docs

* fix unused variable warning

* update README

* fix bug with image injection

During refactoring the read from the file was removed - oops!

* update documentation

Add a note that usage of the memory object after injection will
also be affected.

Author: bashbaug

README.md

@@ -37,6 +37,7 @@ Troubleshooting steps and answers to frequently asked questions can be found [he
 
 Detailed instructions:
 * [How to Inject Modified Programs](docs/injecting_programs.md)
+* [How to Inject Buffers and Images](docs/injecting_buffers_images.md)
 * [How to Dump and Disassemble Kernel ISA](docs/kernel_isa.md)
 * [How to Capture AUB Files](docs/aubcapture.md)
 * [How to Collect MDAPI Performance Metrics](docs/mdapi.md)

docs/controls.md

@@ -517,7 +517,7 @@ This is the list of options that is implicitly passed to CLANG to build a non-Op
 
 This is the list of options that is implicitly passed to CLANG to build an OpenCL 2.0 SPIR-V module.  Any application-provided build options will be appended to these build options.
 
-### Controls for Dumping Buffers and Images
+### Controls for Dumping and Injecting Buffers and Images
 
 ##### `DumpBufferHashes` (bool)
 
@@ -591,6 +591,14 @@ The Intercept Layer for OpenCL Applications will only dump argument values when
 
 The Intercept Layer for OpenCL Applications will only dump kernel arguments when the enqueue counter is less than this value, inclusive.
 
+##### `InjectBuffers` (bool)
+
+If set to a nonzero value, the Intercept Layer for OpenCL Applications will look to inject potentially modified buffer, SVM, and USM contents before calls to clEnqueueNDRangeKernel().  Only buffers that are kernel arguments for the kernel being enqueued may be injected.  The filename to inject will have the form "Enqueue\_\<Enqueue Number\>\_Kernel\_\<Kernel Name\>\_Arg\_\<Argument Number\>\_Buffer\_\<Unique Memory Object Number\>.bin", which matches the filename for dumped buffers.
+
+##### `InjectImages` (bool)
+
+If set to a nonzero value, the Intercept Layer for OpenCL Applications will look to inject potentially modified image contents before calls to clEnqueueNDRangeKernel().  Only images that are kernel arguments for the kernel being enqueued may be injected.  The filename to inject will have the form "Enqueue\_\<Enqueue Number\>\_Kernel\_\<Kernel Name\>\_Arg\_\<Argument Number\>\_Image\_\<Unique Memory Object Number\>\_\<Width\>x\<Height\>x\<Depth\>\_\<Element Size\>bpp.raw", which matches the filename for dumped images.
+
 ### Device Partitioning Controls
 
 ##### `AutoPartitionAllDevices` (bool)

docs/injecting_buffers_images.md

@@ -0,0 +1,70 @@
+# Using the Intercept Layer for OpenCL Applications to Inject Buffers and Images
+
+Buffer and Image Injection allows an application to experiment with different kernel inputs.
+A typical use-case is to dump buffers or images while an application is running on a known good device or driver, and then to inject these known good buffers or images at certain points during application execution on a different device or with a different driver that is producing incorrect results.
+This can help to identify differences that are irrelevant and unrelated to the incorrect results from important differences.
+
+The process to inject buffers and images is very similar to the process to inject programs, and requires several steps.
+
+## Step 1: Dump Programs and Program Options for Injection
+
+Run your application with `DumpBuffersBeforeEnqueue` or `DumpImagesBeforeEnqueue` enabled. This dumps buffers and images that are set as kernel arguments to the directory:
+
+    %SYSTEMDRIVE%\Intel\CLIntercept_Dump\<Process Name>\memDumpPreEnqueue (Windows)
+
+or:
+
+    ~/CLIntercept_Dump/<Process Name>/memDumpPreEnqueue (Linux)
+
+Because this can generate a lot of data, you may want to dump buffers and images only for a small region of the application's execution, using controls like `DumpBuffersMinEnqueue`, `DumpImagesMaxEnqueue`, or `DumpBuffersForKernel`.
+
+**Note.** If the `AppendPid` option is enabled, dumps will be saved to the `<Process Name>.PID` directory, but the `<Process Name>` directory without the process ID will be used for buffer and image injection.
+
+## Step 2: Find the Buffer or Image to Inject
+
+Look through the dump directory to find the buffer or image you would like to inject.
+
+## Step 3: Copy the Buffer or Image
+
+Copy the buffer or image file to inject to the directory:
+
+    %SYSTEMDRIVE%\Intel\CLIntercept_Dump\<Process Name>\Inject (Windows)
+
+or:
+
+    ~/CLIntercept_Dump/<Process Name>/Inject (Linux)
+
+This is the directory that is searched when looking for buffers or images to inject.
+Note that this is the same directory that is used to inject programs, and is a subdirectory of the top-level dump directory.
+
+## Step 4: Modify the Buffer or Image (Optional)
+
+For some use-cases you will want to manually modify the buffer or image before injecting.
+For other use-cases, you will inject the dumped buffer or image, unmodified.
+
+## Step 5: Set the InjectBuffers or InjectImages Controls and Go!
+
+If all goes well you will see a line similar to
+
+    Injecting buffer file: <file name>
+
+or:
+
+    Injecting image file: <file name>
+
+in your log.
+
+## Notes:
+
+* The process to inject buffers also works for SVM or USM allocations.
+* For buffers, the size of the buffer must be at least as big as the file to inject.
+If the size of the buffer is smaller than the size of the file then injection will fail.
+If the size of the buffer is larger than the size of the file then only the initial contents of the buffer will be modified.
+* For OpenCL images, the size of the image must match the size of the file or injection will fail.
+* Because buffer and image injection modifies the memory object itself and not a copy of the memory object, any usage of the memory object after injection will also be affected.
+
+---
+
+\* Other names and brands may be claimed as the property of others.
+
+Copyright (c) 2024, Intel(R) Corporation

intercept/src/controls.h

@@ -118,7 +118,7 @@ CLI_CONTROL( std::string,   SPIRVDis,                               "spirv-dis",
 CLI_CONTROL( std::string,   DefaultOptions,                         "-cc1 -x cl -cl-std=CL1.2 -D__OPENCL_C_VERSION__=120 -D__OPENCL_VERSION__=120 -emit-spirv -triple=spir", "This is the list of options that is implicitly passed to CLANG to build a non-OpenCL 2.0 SPIR-V module.  Any application-provided build options will be appended to these build options." )
 CLI_CONTROL( std::string,   OpenCL2Options,                         "-cc1 -x cl -cl-std=CL2.0 -D__OPENCL_C_VERSION__=200 -D__OPENCL_VERSION__=200 -emit-spirv -triple=spir", "This is the list of options that is implicitly passed to CLANG to build an OpenCL 2.0 SPIR-V module.  Any application-provided build options will be appended to these build options." )
 
-CLI_CONTROL_SEPARATOR( Controls for Dumping Buffers and Images: )
+CLI_CONTROL_SEPARATOR( Controls for Dumping and Injecting Buffers and Images: )
 CLI_CONTROL( bool,          DumpBufferHashes,                       false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will dump hashes of a buffer, SVM, or USM allocation rather than the full contents of the buffer.  This can be useful to identify which kernel enqueues generate different results without requiring a large amount of disk space." )
 CLI_CONTROL( bool,          DumpImageHashes,                        false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will dump hashes of an image rather than the full contents of the image.  This can be useful to identify which kernel enqueues generate different results without requiring a large amount of disk space." )
 CLI_CONTROL( bool,          DumpArgumentsOnSet,                     false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will dump the argument value on calls to clSetKernelArg(). Arguments are dumped as raw binary data.  The filenames will have the form \"SetKernelArg_<Enqueue Number>_Kernel_<Kernel Name>_Arg_<Argument Number>.bin\"." )
@@ -137,6 +137,8 @@ CLI_CONTROL( cl_uint,       DumpImagesMinEnqueue,                   0,     "The
 CLI_CONTROL( cl_uint,       DumpImagesMaxEnqueue,                   UINT_MAX, "The Intercept Layer for OpenCL Applications will only dump image kernel arguments when the enqueue counter is less than this value, inclusive." )
 CLI_CONTROL( cl_uint,       DumpArgumentsOnSetMinEnqueue,           0,     "The Intercept Layer for OpenCL Applications will only dump argument values when the enqueue counter is greater than this value, inclusive." )
 CLI_CONTROL( cl_uint,       DumpArgumentsOnSetMaxEnqueue,           UINT_MAX, "The Intercept Layer for OpenCL Applications will only dump kernel arguments when the enqueue counter is less than this value, inclusive." )
+CLI_CONTROL( bool,          InjectBuffers,                          false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will look to inject potentially modified buffer, SVM, and USM contents before calls to clEnqueueNDRangeKernel().  Only buffers that are kernel arguments for the kernel being enqueued may be injected.  The filename to inject will have the form \"Enqueue_<Enqueue Number>_Kernel_<Kernel Name>_Arg_<Argument Number>_Buffer_<Unique Memory Object Number>.bin\", which matches the filename for dumped buffers." )
+CLI_CONTROL( bool,          InjectImages,                           false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will look to inject potentially modified image contents before calls to clEnqueueNDRangeKernel().  Only images that are kernel arguments for the kernel being enqueued may be injected.  The filename to inject will have the form \"Enqueue_<Enqueue Number>_Kernel_<Kernel Name>_Arg_<Argument Number>_Image_<Unique Memory Object Number>_<Width>x<Height>x<Depth>_<Element Size>bpp.raw\", which matches the filename for dumped images." )
 
 CLI_CONTROL_SEPARATOR( Device Partitioning Controls: )
 CLI_CONTROL( bool,          AutoPartitionAllDevices,                false, "If set to a nonzero value, the Intercept Layer for OpenCL Applications will automatically partition parent devices and return all parent devices and all sub-devices." )

intercept/src/dispatch.cpp

@@ -4795,6 +4795,7 @@ CL_API_ENTRY cl_int CL_API_CALL CLIRN(clEnqueueNDRangeKernel)(
             local_work_size );
         DUMP_BUFFERS_BEFORE_ENQUEUE( kernel, command_queue );
         DUMP_IMAGES_BEFORE_ENQUEUE( kernel, command_queue );
+        INJECT_MEMORY_OBJECTS( kernel, command_queue );
         CHECK_AUBCAPTURE_START_KERNEL(
             kernel,
             work_dim,