DETAILED ACTION
Statement of claims
The present amended application includes:
Claims 11, 9, 15 and 21 were amended.
Claims 1-26 remain pending in the application. Claims 1-26 are being considered on the merits.
Notice of Pre-AIA or AIA Status
The present application, filed on or after March 16, 2013, is being examined under the first inventor to file provisions of the AIA .
Information Disclosure Statement
The information disclosure statement (IDS) submitted on 10/03/2025 and 01/28/2026 . The submission is in compliance with the provisions of 37 CFR 1.97. Accordingly, the information disclosure statement is being considered by the examiner.
Continued Examination Under 37 CFR 1.114
A request for continued examination under 37 CFR 1.114, including the fee set forth in 37 CFR 1.17(e), was filed in this application after final rejection. Since this application is eligible for continued examination under 37 CFR 1.114, and the fee set forth in 37 CFR 1.17(e) has been timely paid, the finality of the previous Office action has been withdrawn pursuant to 37 CFR 1.114. Applicant's submission filed on January 22, 2026 has been entered.
Response to Arguments
Applicant argues that:
“Therefore, neither Gabriel nor Ba, alone or in combination, disclose at least, "changing function code of one or more nodes of the executable graph code, without causing the instantiated executable graph to be re-instantiated, to different function code indicated by a second one or more parameters to the API call," as recited in amended claim 9"
Examiner respectfully disagree and submit that:
Applicant’s arguments with respect to the newly added limitations have been considered but are moot because the arguments do not apply to the reference Venkataramani et al. (US 2018/0157471) and Bo Qiao et al. “The Best of Both Worlds: Combining CUDA Graph with an Image Processing DSL”, Published 2020-07-01. being used in the current rejection
Claim Rejections - 35 USC § 103
The following is a quotation of 35 U.S.C. 103 which forms the basis for all obviousness rejections set forth in this Office action:
A patent for a claimed invention may not be obtained, notwithstanding that the claimed invention is not identically disclosed as set forth in section 102, if the differences between the claimed invention and the prior art are such that the claimed invention as a whole would have been obvious before the effective filing date of the claimed invention to a person having ordinary skill in the art to which the claimed invention pertains. Patentability shall not be negated by the manner in which the invention was made.
Claim(s) 1-6, 8-18, 20-22 and 24-26 are rejected under 35 U.S.C. 103 as being unpatentable over Venkataramani et al. (US 2018/0157471, Venkataramani hereinafter) in view of Bo Qiao et al. “The Best of Both Worlds: Combining CUDA Graph with an Image Processing DSL”, Bo hereinafter, Published 2020-07-01).
Claim Rejections - 35 USC § 103
The following is a quotation of 35 U.S.C. 103 which forms the basis for all obviousness rejections set forth in this Office action:
A patent for a claimed invention may not be obtained, notwithstanding that the claimed invention is not identically disclosed as set forth in section 102, if the differences between the claimed invention and the prior art are such that the claimed invention as a whole would have been obvious before the effective filing date of the claimed invention to a person having ordinary skill in the art to which the claimed invention pertains. Patentability shall not be negated by the manner in which the invention was made.
Claim(s) 1-26 are rejected under 35 U.S.C. 103 as being unpatentable over Venkataramani et al. (US 2018/0157471, Venkataramani hereinafter) in view of Bo Qiao et al. “The Best of Both Worlds: Combining CUDA Graph with an Image Processing DSL”, Bo hereinafter, Published 2020-07-01).
As to claim 1, Venkataramani teaches a non-transitory machine-readable medium having stored thereon instructions that, in response to execution by one or more processors, causes the one or more processors to at least:
in response to an application programming interface (API) call (e.g., para 0049] execution on the host until all issued CUDA calls “, “To synchronize the host with respect to a specific stream, a ‘cudaStreamSynchronize(stream)’ API may be use”), modify executable graph code of an instantiated executable graph indicated by a first one or more parameters of the API call by changing function code of one or more nodes of the executable graph code (e.g., para [0120] To make the stream assignments, the dependency analyzer 418 may create a kernel dependency graph, which may be a directed graph having nodes and edges. The nodes may correspond to the kernels, and a special node may correspond to the CPU. Edges between two kernel nodes may represent a data dependency between the two kernels, for example one kernel may compute a variable that is used by the other kernel. The kernel creation unit 416 may analyze the kernel dependency graph and may apply a partitioning algorithm, such as a clique partitioning algorithm, to the kernel dependency graph to assign each kernel node to a clique. The clique partitioning algorithm may assign kernel nodes to densely connected components, such as cliques, so that the number of edges between cliques is minimized. Each clique may represent a stream, e.g., a CUDA asynchronous stream. The kernel creation unit 416 also may insert ‘cudaDeviceSynchronization’ calls at locations in the generated code 324 that correspond to the edges between cliques of the kernel dependency graph.
para [0122] The kernel creation unit 416 may modify an IR for the source program 312 to include statements that define a kernel”, “statements for launching the kernel”, “ The kernel creation unit 416 may include cudaMalloc APIs for one or more kernel variables and may also include cudaMemcpy APIs for one or more input/output (I/O) kernel variables ” and “ one or more IRs may be graph-based, object-oriented structures”, “ the IRs may be in the form of a hierarchical Data Flow Graph (DFG) and/or a Parallel Intermediate Representation (PIR), which may include a plurality of IR objects, such as nodes, which may represent operators of the source program 312, interconnected by edges, which may represent data flow. The nodes of the PIR may present components corresponding to portions of the source program 312, such as functions and/or operations, and the edges may represent data and/or control flow” in para [0080] ) , to different function code (e.g., see FIG. 7 and 8) indicated by a second one or more parameters to the API call (e.g., see para 60, “ one or more optimizations that modify portions of the IRs of the source program into forms suitable for execution on the device” , “ analyze the use of variables within the source program through examination of one or more of the IRs”, “modify one or more of the IRs of the source program to efficiently utilize the PPU's memory hierarchy”).
Venkataramani Does not teach changing function code without causing the instantiated executable graph to be re-instantiated.
Bo Qiao teaches changing function code (e.g., “performs the source-to source transformation using two internal libraries” in page 4) of an instantiated executable graph (e.g., Graph API”, Figure 1 for “Cuda Runtime” and “II. CUDA GRAPH Utilizing the CUDA graph API follows a three-stage execution model: Graph definition, graph instantiation, and graph execution.” , “The second stage is the graph instantiation. After defining the graph, the whole workflow is known, and analysis can be performed to enable certain optimizations, such as to determine the number of possible parallel streams or to prepare the kernels such that they can be launched as fast as possible during execution.” In page in page 2) without causing the instantiated executable graph to be re- instantiated (e.g., see FIG. 1, “Figure 1: Combining CUDA graph with Hipacc.”, page 2, “we extend the workflow optimizations in CUDA graph with additional techniques such as concurrent kernel execution with complementary resources and kernel fusion” and “A. Framework”, “the Hipacc compiler traverses the AST and performs the source-to source transformation using two internal libraries: Analyzer gathers information for analysis and optimizations, e.g., domain knowledge such as the kernel resource usage, compute pattern, kernel data dependence, or device information such as the compute capability. Based on the obtained domain- and architecture-specific knowledge, optimizations are applied. Hipacc performs optimizations in two steps: First, single kernel optimizations such as memory padding are applied, and efficient CUDA code is generated for each user-defined kernel.”, “Analyzer and Rewriter work iteratively to update the analysis information and implement the desired transformations. In this work, when both kernel fusion and concurrent kernel execution are available, the transformation on concurrent execution takes precedence over fusion, because we think a better resource utilization has higher priority. In the end, the backend of Hipacc generates a CUDA graph implementation with reduced launch latency and a fully optimized workflow” in page 4 and page 5 .
According to applicant’s specification in para [0053] …in a task graph, a workload that includes multiple tasks is organized as a directed graph, where each node corresponds to a task to be performed, and each directed edge between two nodes corresponds to a data dependency, an execution dependency, or some other dependency between two nodes.
….an executable graph may be generated from a task graph using instantiation. In at least one embodiment, because converting a task graph to an executable graph has access to an entire task graph, various optimizations may be performed, which may reduce an overall execution time of a workload. In at least one embodiment, an executable graph may be used multiple times to have computing resources perform a same workload without having to be regenerated from a task graph. In at least one embodiment, tasks include kernel functions described by function code to be performed by one or more GPUs.
Thus, the “Rewriter work iteratively to update the analysis information and implement the desired transformations”, “the transformation on concurrent execution takes precedence over fusion”, therefore, changing function code of executable graph code the instantiated executable graph without causing the instantiated executable graph to be re- instantiated). Thus, it would have been obvious to one of ordinary skill in the art before the effective filing date of the claimed invention to modify the method of Venkataramani by adopting the teachings of Bo to “improve the ability of CUDA graph to exploit workflow optimizations, e.g., concurrent kernel executions with complementary resource occupancy” (see Abstract, Bo).
As to claim 2, Venkataramani teaches wherein the one or more nodes of the executable graph code are indicated to the API (e.g., para 80, “The nodes of the PIR may present components corresponding to portions of the source program 312, such as functions and/or operations”, and para [0122] The kernel creation unit 416 may include cudaMalloc APIs for one or more kernel variables and may also include cudaMemcpy APIs”) .
As to claim 3, Venkataramani teaches wherein the instructions, which if performed by the one or more processors, are to cause the one or more processors to replace code for a first graphics processing unit (GPU) kernel with code for a second GPU kernel in the executable graph code (e.g., para 60, “identify portions of an IR of the source program that can be replaced with calls to a predefined PPU-optimized library” for “a kernel dependency graph” and “kernels, that are concurrently executed, “executed in parallel, by the GPU” in para 41, 120 ).
As to claim 4, Venkataramani teaches wherein the instructions, which if performed by the one or more processors, are to cause the one or more processors to modify the executable graph code based, at least in part, on a parameter to the API that identifies function code to be used (e.g., para 60, “apply one or more optimizations that modify portions of the IRs of the source program into forms suitable for execution on the device”, “the use of variables within the source program through examination of one or more of the IRs, and may determine a reduced set of memory transfers between the host and device. The systems and methods also may modify one or more of the IRs of the source program to efficiently utilize the PPU's memory hierarchy”).
As to claim 5, Venkataramani teaches wherein the API is to modify a single node in the executable graph code, and the instructions, which if performed by the one or more processors. are to cause the one or more processors to change a first CPU kernel to a second CPU kernel to be performed by the single node (e.g., para [0086] The partitioning engine 414 may partition one or more of the IRs into segments for host (CPU) and device (PPU) execution, as indicated at step 516. For example, the regions identified as suitable for parallel execution may be partitioned into the segments to be executed by the device (PPU), while other regions may be partitioned for execution by the host (CPU). The kernel creation unit 416 may create kernels within one or more IRs of the source program for the segments to be executed by the device (PPU), as indicated at step 518. For example, the kernel creation unit 416 may modify the one or more IRs of source program 312 by adding CUDA or OpenCL declarations that define and launch kernels).
As to claim 6 , Venkataramani teaches wherein the API is to modify the executable graph code based, at least in part, on changing function code to be performed by a graphics processing unit (CPU), and the instructions, which if performed by the one or more processors, are to cause the one or more processors to update information to be copied to the CPU based, at least in part, on the modified executable graph code (e.g., para 196, “an index of maximum magnitude portion that may be replaced with a ‘cublasiSamax’ function from the CUDA cuBLAS function library, a Rank-1 Update portion that may be replaced with a ‘cublasZger’ function, a ‘cublasZgerc’ function, or a ‘cublasZgeru’ function from the CUDA cuBLAS function library, and a QR factorization portion that may be replaced with a ‘cusolverDunormqr’ function from the cuSOLVER function library or a ‘magma_Sgeq3_gpu’ function from the MAGMA function library” and “ a modified version of the source program 312, for example modified to include kernel creation and kernel launch statements, among other changes” in para 72).
As to claim 8, Venkataramani teaches wherein the instructions, which if performed by the one or more processors, are to cause the one or more processors to change function code of one or more nodes of the executable graph code and to change one or more parameters to be used by the function code (e.g., para 60, “apply one or more optimizations that modify portions of the IRs of the source program into forms suitable for execution on the device”, “the use of variables within the source program through examination of one or more of the IRs, and may determine a reduced set of memory transfers between the host and device”).
As to claim 9-14 see rejection of claim1-3, 5 and 8 above. Venkataramani teaches further one or more processors, comprising: circuitry ( see FIG. 2).
As to claims 15-18, 20, see rejection of claims 1-6 and 8 above. Venkataramani teaches further A system, comprising : one or more processors and one or more memories to store the executable graph code ( see FIG. 3)
As to claims 21-24, see rejection of claims 1-6 and 8 above
As to claim 25, Venkataramani teaches wherein modifying executable graph code includes changing function code of a node to be specified by a parameter to the API, and the method includes selecting a buffer for kernel data based, at least in part, on the changed function code (e.g., para [0258] then the kernel creation unit 416 may create a kernel with grid dimension, number of blocks, and number of threads computed from the parameters, e.g., number of iterations” ).
As to claim 26, Venkataramani teaches wherein the executable graph code is to be specified by a first parameter to the API, a node of the executable graph code is to be specified by a second parameter to the API, function code to be performed by one or more graphics processing units (GPUs) is to be specified based, at least in part, on a third parameter to the API, and modifying executable graph code includes changing the node to perform the function code (e.g., para [0258] then the kernel creation unit 416 may create a kernel with grid dimension, number of blocks, and number of threads computed from the parameters, e.g., number of iterations”. Thus, another one of the parameters include third parameter to the API ).
Claim(s) 7,19 and 23 are rejected under 35 U.S.C. 103 as being unpatentable over Venkataramani et al. (US 2018/0157471, Venkataramani hereinafter) in view of Bo Qiao et al. “The Best of Both Worlds: Combining CUDA Graph with an Image Processing DSL”, Bo hereinafter, Published 2020-07-01), as applied to claims 1, 15 and 21 above, and further in view of Guerfinkel et al. (US 2021/0149734, Gurfinkel hereinafter).
As to claim 7, Gurfinkel teaches wherein the instructions, which if performed by the one or more processors, are to cause the one or more processors to change a node of the executable graph code to be a null node (e.g., para [0291] In at least one embodiment, as each task is added as a node to a task graph, a creation order of each node is associated with each node. In at least one embodiment, a creation order associated with each node corresponds to an ordinal number indicating whether each node is a first, second, third, or so forth node created and added to a task graph. In at least one embodiment, a creation order of each node may correspond to a count value of a counter that is initialized (e.g., to zero or one) when a task graph is initially created and is incremented each time a new node is added to a task graph.) . Thus, it would have been obvious to one of ordinary skill in the art before the effective filing date of the claimed invention to modify the method of Venkataramani and Bo by adopting the teachings of Gurfinkel to “ reduce an overall execution time of a workload.” (see para 287, Gurfinkel).
As to claims 19 and 23 see rejection of claim 7 above.
Conclusion
Any inquiry concerning this communication or earlier communications from the examiner should be directed to ABDOU K SEYE whose telephone number is (571)270-1062. The examiner can normally be reached M-F 9-5:30.
Examiner interviews are available via telephone, in-person, and video conferencing using a USPTO supplied web-based collaboration tool. To schedule an interview, applicant is encouraged to use the USPTO Automated Interview Request (AIR) at http://www.uspto.gov/interviewpractice.
If attempts to reach the examiner by telephone are unsuccessful, the examiner’s supervisor, Pierre Vital can be reached at 5712724215. The fax phone number for the organization where this application or proceeding is assigned is 571-273-8300.
Information regarding the status of published or unpublished applications may be obtained from Patent Center. Unpublished application information in Patent Center is available to registered users. To file and manage patent submissions in Patent Center, visit: https://patentcenter.uspto.gov. Visit https://www.uspto.gov/patents/apply/patent-center for more information about Patent Center and https://www.uspto.gov/patents/docx for information about filing in DOCX format. For additional questions, contact the Electronic Business Center (EBC) at 866-217-9197 (toll-free). If you would like assistance from a USPTO Customer Service Representative, call 800-786-9199 (IN USA OR CANADA) or 571-272-1000.
/ABDOU K SEYE/ Examiner, Art Unit 2198
/TUAN C DAO/ Primary Examiner, Art Unit 2198
Prosecution Insights