A GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation

Abstract

The present invention discloses a method for testing the multi-level capacity and extreme capacity of a GPU cache based on bidirectional approximation. Based on the influence of GPU caches at various levels on a computing process, a capacity test kernel function is constructed. The occupancy of a GPU cache to be tested is changed by adjusting the number of loops in the capacity test kernel function and the value of an input array length. The limit capacity of the GPU to be tested is obtained through a forward approximation test. The result obtained through a backward approximation test is then combined to determine the number of cache levels and the capacity of caches at various levels contained in the GPU to be tested, thereby achieving accurate measurement of the limit capacity of the GPU cache and the capacity of caches at various levels.

Description

Technical Field

[0001] The invention belongs to the technical field of graphics testing, and in particular relates to a method for testing multi-level capacity and extreme value capacity of a GPU cache based on bidirectional approximation. Background Art

[0002] The general-purpose computing GPU architecture is built around an extended array of streaming multiprocessors (SMs), which are replicated to achieve GPU hardware parallelism. Typically, a GPU contains two levels of cache, where each SM has a first-level cache, and all SMs share a second-level cache, which is then connected to global registers via the second-level cache. However, there are also special cases where a GPU contains a third-level cache.

[0003] Existing general computing methods such as OpenCL and CUDA provide special APIs to obtain the cache capacity in the GPU. However, the cache capacity obtained through the function interface is a fixed value preset in the driver, not the actual value obtained through measurement. In addition, the throughput-based method can only obtain the amount of data accessed per unit time in the general computing GPU cache, rather than the actual capacity of the cache. Summary of the invention

[0004] In view of this, the present invention provides a GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation, which realizes accurate measurement of the ultimate capacity of the GPU cache and the actual values ​​of the cache capacity at each level.

[0005] The present invention provides a method for testing multi-level capacity and extreme capacity of a GPU cache based on bidirectional approximation, comprising the following steps:

[0006] Step 1: Obtain the maximum value of the nominal cache of the GPU to be tested as the first capacity C, and obtain the number of parallel threads TN supported by the GPU to be tested; select a grayscale image as the image to be tested required for the test, and record the total number of pixels contained in the image to be tested as the total number of calculated points PN;

[0007] Step 2, construct a capacity test kernel function for testing the GPU cache to be tested, the capacity test kernel function executes a loop operation of loop number M, the loop body is to first assign the value of the first element of the input array to the intermediate variable, and then perform a cyclic left shift operation on the input array, wherein the length of the input array is the first length N, and the elements in the input array and the intermediate variables are all byte variables; the calculation amount of the capacity test kernel function is M*(N+1), and the execution time of moving an element by one bit is T;

[0008] Step 3: Set (a*C-1) as the initial value of N, a as the coefficient, set a=1, and then set the value of M that satisfies the condition [M*(N+1)*T*PN] / TN which is much larger than the time consumed by the GPU to read the image to be tested as the initial value of M. 0 ; According to the set initial values ​​of M and N, create multiple test threads, each test thread executes the capacity test kernel function respectively, and the pixel point in the tested image is used as the first element of the input array in the test thread;

[0009] Step 4: Start all test threads to perform forward approximation test. When the tested image is processed, terminate all test threads and save the (N+1), test execution time and output result used in this execution as the forward approximation test result. If the difference between the output result and the tested image is less than the threshold, set a=2*a and N=(a*C-1) to update the value of N, and the value of M to M / 2, and execute step 4. Otherwise, save the current values ​​of a, N and M, take the current (N+1) as the extreme value capacity, and execute step 5.

[0010] Step 5, let a=a / 2 and N=(a*C-1) update the value of N, and the value of M is M*2; start all test threads to perform the backward approximation test, and terminate all test threads when the tested image is processed, and save the (N+1), test execution time and output results used in this execution as the backward approximation test result; if a is not less than 1, execute step 5, otherwise execute step 6;

[0011] Step 6: Establish a forward cache capacity and computation time curve according to the forward approximation test result, establish a backward cache capacity-computation time curve according to the backward approximation test result, select multiple pairs of curves whose similarity between the forward cache capacity-computation time curve and the backward cache capacity-computation time curve is greater than a threshold, calculate the mutation points on the curves, and take the mutation points where the subsequent computation time increases and remains unchanged within a certain period of time as valid mutation points. The number of valid mutation points is the number of levels of the GPU cache to be tested, and the cache capacity corresponding to the valid mutation points is the capacity of the cache at each level.

[0012] Furthermore, in step 3, the time consumed by the GPU to be tested to read the image to be tested is obtained by recording the time required to perform multiple reading operations and then taking an average value.

[0013] Furthermore, the method for obtaining the number of parallel threads TN supported by the GPU to be tested in step 1 is: obtaining it in a block manner using a two-dimensional kernel function for image data.

[0014] Furthermore, the image to be measured is composed of two grayscale images with differences alternately.

[0015] Furthermore, the images under test are two images, Baboon and Goldhill.

[0016] Furthermore, after the capacity test kernel function is executed for the first time, the binary code is saved, and the capacity test kernel function is updated by modifying the number of loops and the first length in the binary code.

[0017] Furthermore, the forward cache capacity-computation time curve described in step 6 is established according to the forward approximation test results, and the backward cache capacity-computation time curve is established according to the backward approximation test results. The implementation method is: by sampling the forward approximation test results and the backward approximation test results, and using the sampled data to establish the curve.

[0018] Furthermore, the sampling method is: first determine the preliminary sampling times k, measure k groups of data, then calculate the initial sampling mean u' and initial sampling variance σ^', and use the exponential distribution function to determine the qualified sampling times under the conditions of confidence α and error e.

[0019] Beneficial effects:

[0020] The present invention constructs a capacity test kernel function based on the influence of GPU caches at all levels on the computing process, changes the occupancy of the GPU cache to be tested by adjusting the number of loops in the capacity test kernel function and the value of the input array length, obtains the limit capacity of the GPU to be tested through a forward approximation test, and then determines the number of cache levels and the capacity of caches at all levels contained in the GPU to be tested by combining the results obtained by the backward approximation test, thereby achieving accurate measurement of the limit capacity of the GPU cache and the capacity of caches at all levels. BRIEF DESCRIPTION OF THE DRAWINGS

[0021] Figure 1 This is a schematic diagram of the effect of the number of threads on the computation time in the test process involved in the present invention.

[0022] Figure 2 A schematic diagram of the calculation process of a capacity test kernel function constructed in a GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation provided by the present invention.

[0023] Figure 3 A schematic diagram of a cache capacity-computation time curve obtained by a GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation provided by the present invention. DETAILED DESCRIPTION

[0024] The present invention is described in detail with reference to the following embodiments.

[0025] According to the structure of general-purpose computing GPUs, the capacity and processing speed of each level of cache are different. The first-level cache (L1 cache) has a small capacity and a fast speed, while the second-level cache (L2 cache) has a large capacity and a slow speed. If all calculations of the computing core are implemented on the first-level cache, the fastest computing speed will be achieved. If the required storage capacity exceeds the capacity of the first-level cache and the second-level cache is required, more computing time will be required to maintain the same amount of computing.

[0026] The present invention provides a method for testing the multi-level capacity and extreme capacity of a GPU cache based on bidirectional approximation. The core idea is to construct a capacity test kernel function based on the influence of GPU caches at various levels on the computing process, change the occupancy of the GPU cache to be tested by adjusting the number of loops in the capacity test kernel function and the value of the input array length, obtain the limit capacity of the GPU to be tested through a forward approximation test, and then determine the number of cache levels contained in the GPU to be tested and the capacity of each cache level by combining the results obtained by a backward approximation test.

[0027] The present invention provides a method for testing multi-level capacity and extreme capacity of a GPU cache based on bidirectional approximation, which specifically includes the following steps:

[0028] Step 1: Obtain the multi-level capacity of the GPU cache to be tested, which is marked in the driver program, and take the maximum value of the multi-level capacity as the first capacity C, and obtain the number of parallel threads TN supported by the GPU to be tested; select a grayscale image as the image to be tested required for the test, and record the total number of pixels contained in the image to be tested as the total number of calculated points PN. The multi-level capacity of the GPU cache to be tested, which is marked in the driver program, is usually the value pre-written by the manufacturer.

[0029] Whether the number of threads used in the test process matches the support of the GPU to be tested will affect the execution time of the test process. The impact of the number of threads used on the calculation time is as follows: Figure 1 As shown, the present invention needs to set the test process according to the optimal number of threads supported by the GPU to be tested. The number of parallel threads TN supported by the GPU to be tested can be obtained by: using a two-dimensional kernel function block method for image data to obtain the optimal number of threads, dx*dy=block thread number (a multiple of 2), dx is the number of rows in the two-dimensional space, dy is the number of columns in the two-dimensional space, and as the exploration proceeds, the measurement curve moves downward as a whole. After reaching the optimal value, the measurement curve will not shift significantly if the number of block threads is increased.

[0030] Taking OpenCL as an example, set the size of the block calculation to size_t local_size[2]={dx,dy};

[0031] Serial number dx value dy value Number of blocking threads 1. 2 2 4 2. 2 4 8 3. 4 4 16 4. 4 8 32 5. 8 8 64 6. 8 16 128 7. 16 16 256 8. 16 32 512 9. 32 32 1024

[0032] Limited by the chip size, the optimal number of block threads is often less than 1024, so this can be used as the stopping threshold.

[0033] Call the kernel function for calculation, specifically clEnqueueNDRangeKernel (command queue, kernel function name, dimension parameter set to offset, input data size, i.e. image size, block calculation size, and other parameters), gradually reduce the number of block threads to make the measurement curve move downward as a whole, and after reaching the optimal value, increase the number of block threads so that the measurement curve will not shift significantly, and the optimal number of parallel threads can be obtained.

[0034] In order to better meet the requirements of scalar calculation, the present invention selects a grayscale image as the image to be measured, and each pixel in the grayscale image is a single byte.

[0035] In order to avoid repeated input of the same tested image, which would make it difficult to identify that the cached content has not been updated, and to further improve the accuracy of the test, the present invention selects two different grayscale images, namely image A and image B, and alternately uses them as the tested images. For example, two grayscale image test standard images Baboon and Goldhill are selected from the standard image library, and the image size is 512*512.

[0036] The difference between the two images is calculated according to the cross entropy calculation formula of the two images. The calculation formula is as follows:

[0037]

[0038] Among them, p(i) is the histogram statistic of the pixel with grayscale i in image A, that is, p(i) = the statistical number of pixels with grayscale i in the image / the total number of pixels in the image; q(i) is the histogram statistic of the pixel with grayscale i in image B, and the range of i is [0,255].

[0039] Since the cross entropy between the two grayscale images being tested is a constant, the cross entropy between the output images obtained after the two grayscale images are processed by the capacity test kernel function constructed by the present invention is correct only when the output image is equal to the constant or the difference with the constant is less than the set threshold. At this time, the test result data obtained is valid, and the rest of the data can be discarded, which can effectively avoid the output image of the processing process constructed for executing the present invention from interfering with the statistics of the calculation time, and can ensure that the calculation cache can be updated each time.

[0040] Step 2: Construct a capacity test kernel function for testing the GPU cache to be tested. The calculation process of the capacity test kernel function is as follows: Figure 2As shown, the capacity test kernel function executes a loop operation with a set number of loops M. The loop body of the loop operation is to first assign the value of the first element of the input array to the intermediate variable, and then perform a circular left shift operation on the input array, wherein the length of the input array is the first length N, and the elements in the input array and the intermediate variables are all byte variables; the computational complexity of the capacity test kernel function is M*(N+1), and the execution time of moving an element in the input array by one position is T.

[0041] The purpose of the capacity test kernel function constructed by the present invention is to obtain the final output result by performing a cyclic left shift operation on the input array in the loop body. If the execution process of the loop body is complete and normal, the final output result is the same as the initial value of the input array. Otherwise, there will be a large difference between the final output result and the initial value of the input array, so as to judge the correctness of the execution process.

[0042] In order to further improve the testing efficiency, the present invention saves the binary code compiled and generated by the capacity test kernel function after the first execution of the capacity test kernel function. The capacity test kernel function can be updated by modifying the number of loops and the first length in the binary code. The capacity test kernel function can be executed again by only loading the binary code, thereby avoiding repeated compilation of the capacity test kernel function and saving overall testing time.

[0043] Step 3: Set (a*C-1) as the initial value of N, a as the coefficient, and set a=1. Then set the value of M that satisfies the execution time of the capacity test kernel function to be much longer than the time consumed by the GPU under test to read the tested image as the initial value of M. 0 , where the execution time of the capacity test kernel function is [M*(N+1)*T*PN] / TN. The time consumed by the tested GPU to read the tested image can be obtained by taking the average of the time required to perform multiple read operations. According to the set initial values ​​of M and N, multiple test threads are created, each of which is used to execute the capacity test kernel function. The pixel points in the tested image are used as the first element of the input array in the test thread.

[0044] Step 4: Take the image under test as input, start all test threads to perform forward approximation test, start recording test execution time, stop the execution process of all test threads when the image under test is processed, end recording test time, save (N+1), test execution time and output result used in this execution as forward approximation test result. If the difference between the output result and the image under test is less than the threshold, let a=2*a and N=(a*C-1) update the value of N, the value of M is M / 2, and then execute step 4; otherwise, save the current values ​​of a, N and M, take the current (N+1) as the extreme capacity, and execute step 5.

[0045] During the forward approximation test, it is possible that when the storage size occupied by the capacity test kernel function exceeds the capacity limit set by the driver, it can be detected that the test execution time will suddenly drop. At this time, the GPU under test does not actually execute the test process but only executes the input and output operations of the image under test. Therefore, the execution time at this time is only the system I / O time, and the calculated output data does not meet the expected output results. At this time, the cache capacity set by the test is the extreme capacity.

[0046] Step 5, let a=a / 2 and N=(a*C-1) update the value of N, and the value of M is M*2; use the image under test as input, start all test threads to perform the backward approximation test, start recording the test execution time, stop the execution process of all test threads when the image under test is processed, end the recording of the test time, save the (N+1), test execution time and output results used in this execution as the backward approximation test result. If a is not less than 1, execute step 5, otherwise execute step 6.

[0047] Step 6: Establish a forward cache capacity-computation time curve based on the forward approximation test results, and establish a backward cache capacity-computation time curve based on the backward approximation test results. Figure 3 As shown, the similarity between the forward cache capacity-computation time curve and the backward cache capacity-computation time curve is evaluated, and multiple pairs of curves with similarities greater than a threshold are selected. The mutation points on these curves are calculated, and the mutation points where the subsequent calculation time increases and remains unchanged within a certain period of time are taken as valid mutation points. The number of valid mutation points is the number of levels of the GPU cache to be tested, and the cache capacity corresponding to the valid mutation points is the capacity of the cache at each level.

[0048] In order to improve the efficiency of generating curves, the present invention samples the test results. The specific process is: first determine the initial sampling times k, measure k groups of data, and then calculate the initial sampling mean u' and the initial sampling variance σ^'. According to the principle that the calculation time obeys the exponential distribution, the exponential distribution function is used to determine the sampling times that meet the conditions under the conditions of confidence α and error e.

[0049] Example:

[0050] In this embodiment, a GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation provided by the present invention is adopted to measure the GPU cache capacity based on OpenCL and CUDA, which specifically includes the following steps:

[0051] S1. Using the API in general computing methods such as OpenCL and CUDA to return the capacity of the secondary cache of the GPU to be tested as 256K, a capacity test kernel function for testing the cache of the GPU to be tested is constructed. The capacity test kernel function implements the internal loop calculation, specifically including:

[0052] S1.1. Define the input array byte A[N], where the initial value of N is 256*1024-1=262144-1; define the intermediate variable byte B, let B=A[0].

[0053] S1.2. Let m=0, where m is an integer not greater than M.

[0054] S1.3. Let n=0, where n is an integer not greater than N-2.

[0055] S1.4. Let A[n] = A[n+1]. If n is not greater than N-2, then n is incremented by 1 and S1.4 is executed. Otherwise, S1.5 is executed.

[0056] S1.5. At this time, the values ​​of the elements in the input array byte A[N] are A[0]=A[1], A[1]=A[2], ..., A[n-2]=A[n-1], and then A[n-1]=B, that is, A[n-1]=A[0]. If m is not greater than M, m is incremented by 1 and S1.3 is executed. Otherwise, the execution of this capacity test kernel function is terminated.

[0057] S2. Use the grayscale image test standard images Baboon and Goldhill alternately as the tested images, and use the value of M that satisfies the execution time of the capacity test kernel function much longer than the time consumed by the tested GPU to read the tested image as the initial value M of M. 0 .

[0058] To ensure that the total amount of computation in each round of testing remains unchanged, N is increased by approximately double the value in the forward approximation test and M is attenuated by half of the previous value. Conversely, N is attenuated by half of the previous value in the backward approximation test and M is increased by approximately double the value to keep the computation amount M*(N+1) basically unchanged. For example, in the forward approximation test, assuming that the initial value of N is (256*1024-1) and the value of M is 4, then the value of N in the next round is (512*1024-1) and the value of M is 2.

[0059] S3, forward approximation test determines the upper limit capacity of the cache and whether there is a third-level cache (L3 cache). Due to driver limitations, when the cache capacity of the GPU to be tested is set too large during the test, the calculation time will be reduced to the minimum, and at this time, the output result is significantly different from the image under test, and the forward approximation test is terminated. If the GPU under test has a third-level cache, the curve will show a third mutation. If the cache capacity of the test setting is increased after the third mutation, the upper limit capacity will be reached.

[0060] Taking the limit capacity as the starting value, a backward approximation test is performed to obtain a backward approximation test result.

[0061] S4. According to the set sampling method, a forward cache capacity-computation time curve is established based on the forward approximation test results, and a backward cache capacity-computation time curve is established based on the backward approximation test results. The mean square error (MSE) algorithm is used to calculate the mean square error of the two curves, that is, the difference between each corresponding point of the two curves is squared, and then the average value is calculated. The smaller the MSE, the more similar the two curves are. According to the 3 times mean square error calculation formula, for a data set with multiple observation samples, a range is obtained by calculating the variance and multiplying it by 3, which contains approximately 99.7% of the data. The curves that fall within the range of 3 times the mean square error are selected for calculation. When the curve has a clear step, the point is used as the average value of the first-level cache capacity and the second-level cache capacity. The remaining curves are discarded as errors because their correlation is too weak. At the same time, it is necessary to ensure that the number of retained curves is not less than the number of sampling times required to meet the cache test accuracy requirements.

[0062] Serial number Start N Cycle times Calculation times The kernel function accounts for 1 1M-1 1 1M 1M 2 512K-1 2 1M 512K 3 256K-1 4 1M 256K 4 128K-1 8 1M 128K 5 64K-1 16 1M 64K 6 32K-1 32 1M 32K 7 16K-1 64 1M 16K 8 8K-1 128 1M 8K 9 4K-1 256 1M 4K 10 2K-1 512 1M 2K 11 1024-1 1K 1M 1024 12 512-1 2K 1M 512 13 256-1 4K 1M 256 14 128-1 8K 1M 128 15 64-1 16K 1M 64 16 32-1 32K 1M 32 17 16-1 64K 1M 16 18 8-1 128K 1M 8

[0063] The calculation time of the kernel function can be obtained through the special functions of OpenCL or CUDA. CUDA can obtain the execution time of each part through the API, OpenCL can obtain the execution time of the entire kernel function through the API, and can also insert timing points in the main function to obtain it through the timing point difference.

[0064] In summary, the above are only preferred embodiments of the present invention and are not intended to limit the protection scope of the present invention. Any modifications, equivalent substitutions, improvements, etc. made within the spirit and principles of the present invention should be included in the protection scope of the present invention.

Claims

1. A GPU cache multi-level capacity and extreme capacity testing method based on bidirectional approximation, characterized in that: The following steps are involved: Step 1: Obtain the maximum value of the nominal cache of the GPU to be tested as the first capacity C, and obtain the number of parallel threads TN supported by the GPU to be tested; select a grayscale image as the image to be tested required for the test, and record the total number of pixels contained in the image to be tested as the total number of calculated points PN; Step 2, construct a capacity test kernel function for testing the GPU cache to be tested, the capacity test kernel function executes a loop operation of loop number M, the loop body is to first assign the value of the first element of the input array to the intermediate variable, and then perform a cyclic left shift operation on the input array, wherein the length of the input array is the first length N, and the elements in the input array and the intermediate variables are all byte variables; the calculation amount of the capacity test kernel function is M*(N+1), and the execution time of moving an element by one bit is T; Step 3, set (a*C-1) as the initial value of N, a as a coefficient and set a=1, and then set the value of M that satisfies the condition [M*(N+1)*T*PN] / TN which is much larger than the time consumed by the GPU to be tested to read the tested image as the initial value M0 of M; according to the set initial values ​​of M and N, create multiple test threads, each of which executes the capacity test kernel function respectively, and the pixel point in the tested image is used as the first element of the input array in the test thread; Step 4: Start all test threads to perform forward approximation test. When the tested image is processed, terminate all test threads and save the (N+1), test execution time and output result used in this execution as the forward approximation test result. If the difference between the output result and the tested image is less than the threshold, set a=2*a and N=(a*C-1) to update the value of N, and the value of M is M / 2, and execute step 4. Otherwise, save the current values ​​of a, N and M, take the current (N+1) as the extreme value capacity, and execute step 5. Step 5, let a=a / 2 and N=(a*C-1), update the value of N, and the value of M is M*2; start all test threads to perform the backward approximation test, and terminate all test threads when the tested image is processed, and save the (N+1), test execution time and output results used in this execution as the backward approximation test result; if a is not less than 1, execute step 5, otherwise execute step 6; Step 6: Establish a forward cache capacity and computation time curve according to the forward approximation test result, establish a backward cache capacity-computation time curve according to the backward approximation test result, select multiple pairs of curves whose similarity between the forward cache capacity-computation time curve and the backward cache capacity-computation time curve is greater than a threshold, calculate the mutation points on the curves, and take the mutation points where the subsequent computation time increases and remains unchanged within a certain period of time as valid mutation points. The number of valid mutation points is the number of levels of the GPU cache to be tested, and the cache capacity corresponding to the valid mutation points is the capacity of the cache at each level.

2. The GPU cache multi-level capacity and extreme capacity testing method according to claim 1, characterized in that: The time consumed by the GPU under test to read the image under test in step 3 is obtained by recording the time required to perform multiple reading operations and then taking the average value.

3. The GPU cache multi-level capacity and extreme capacity testing method according to claim 1, characterized in that: The method for obtaining the number of parallel threads TN supported by the GPU to be tested in step 1 is: obtaining it by using a two-dimensional kernel function block method for image data.

4. The GPU cache multi-level capacity and extreme capacity testing method according to claim 1, characterized in that: The measured image is composed of two grayscale images with differences alternately.

5. The GPU cache multi-level capacity and extreme capacity testing method according to claim 4, characterized in that: The images under test are two images: Baboon and Goldhill.

6. The GPU cache multi-level capacity and extreme capacity testing method according to claim 1, characterized in that: After the capacity test kernel function is executed for the first time, the binary code is saved, and the capacity test kernel function is updated by modifying the number of loops and the first length in the binary code.

7. The GPU cache multi-level capacity and extreme capacity testing method according to claim 1, characterized in that: The step 6 describes establishing a forward cache capacity-computation time curve based on the forward approximation test results, and establishing a backward cache capacity-computation time curve based on the backward approximation test results. The implementation method is: by sampling the forward approximation test results and the backward approximation test results, and using the sampled data to establish the curve.

8. The GPU cache multi-level capacity and extreme capacity testing method according to claim 7, characterized in that: The sampling method is as follows: first determine the initial sampling times k, measure k groups of data, then calculate the initial sampling mean u' and initial sampling variance σ^', and use the exponential distribution function to determine the sampling times that meet the conditions under the conditions of confidence α and error e.

Images