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 statements (IDS) submitted on 10/30/2025, 03/03/2026 and 04/02/2026. The submission is in compliance with the provisions of 37 CFR 1.97. Accordingly, the information disclosure statements are being considered by the examiner.
Specification
The disclosure is objected to because of the following informalities:
The identification of the incorporated application is incomplete and does not identify the documents being incorporated.
Response to Arguments
Applicant's arguments filed 03/02/2026 have been fully considered but they are not persuasive.
On page 9 applicant argues that Harris only teaches “a single group of thread blocks” and does not teach “a cluster comprising a plurality of groups of one or more blocks of threads”. Examiner respectfully disagrees. On page 2, Harris teaches that Cooperative extends CUDA to “dynamically organize groups of threads” and describes synchronization also on page 2 “both within and across CUDA threads blocks”, Harris further teaches APIs for “ defining, partitioning, and synchronizing groups of threads” and host side APIs for grids whose threads execute concurrently “ to enable synchronization across thread blocks.”
On page 9 applicant argues that a single group is not the same as the claimed plurality of groups. Examiner respectfully disagrees. On page 3, Harris teaches “group partitioning operations”, “Intrinsic groups defined by the CUDA launch API (e.g., thread blocks)” , “A group barrier synchronization operation” and on page 8, Partitioning Groups section “tiled_partition()” partitions a thread block into multiple tiles i.e. multiple groups. Therefore under BRI Harris teaches a cooperative execution having a plurality of groups associated with thread blocks
On page 9 applicant argues that Harris and Nickolls do not teach causing the threads "stop at least until all threads within each of the plurality of groups of the cluster have performed the corresponding barrier instructions." Examiner respectfully disagrees. On page 5, Harris teaches “sync()” , “These perform barrier synchronization among all threads in the group.” Thus, Harris teaches the API/group synchronization framework and Nickolls also teaches stop wait behaviour col 23, line 9-21 “Wait/go registers 1138 keep track of which thread groups have reached which barrier points and are waiting for synchronization to be achieved. In one embodiment, wait/go registers 1108 include a wait/go bit and a barrier identifier BarID field for each of the G SIMD groups that can concurrently execute in core 310. The wait/go bit is set to the wait state (e.g., logic high) when the corresponding SIMD group is waiting at one of the barrier points to synchronize with one or more other SIMD groups and to the go state (e.g., logic low) when the corresponding SIMD group is not waiting at any barrier point; for each group whose wait/go bit is in the wait state, the BarID field is populated with the barrier identifier of the barrier point at which the group is waiting.” See also col 23, line 33 – col 24, line 8.
Double Patenting
The nonstatutory double patenting rejection is based on a judicially created doctrine grounded in public policy (a policy reflected in the statute) so as to prevent the unjustified or improper timewise extension of the “right to exclude” granted by a patent and to prevent possible harassment by multiple assignees. A nonstatutory double patenting rejection is appropriate where the conflicting claims are not identical, but at least one examined application claim is not patentably distinct from the reference claim(s) because the examined application claim is either anticipated by, or would have been obvious over, the reference claim(s). See, e.g., In re Berg, 140 F.3d 1428, 46 USPQ2d 1226 (Fed. Cir. 1998); In re Goodman, 11 F.3d 1046, 29 USPQ2d 2010 (Fed. Cir. 1993); In re Longi, 759 F.2d 887, 225 USPQ 645 (Fed. Cir. 1985); In re Van Ornum, 686 F.2d 937, 214 USPQ 761 (CCPA 1982); In re Vogel, 422 F.2d 438, 164 USPQ 619 (CCPA 1970); In re Thorington, 418 F.2d 528, 163 USPQ 644 (CCPA 1969).
A timely filed terminal disclaimer in compliance with 37 CFR 1.321(c) or 1.321(d) may be used to overcome an actual or provisional rejection based on nonstatutory double patenting provided the reference application or patent either is shown to be commonly owned with the examined application, or claims an invention made as a result of activities undertaken within the scope of a joint research agreement. See MPEP § 717.02 for applications subject to examination under the first inventor to file provisions of the AIA as explained in MPEP § 2159. See MPEP § 2146 et seq. for applications not subject to examination under the first inventor to file provisions of the AIA . A terminal disclaimer must be signed in compliance with 37 CFR 1.321(b).
The filing of a terminal disclaimer by itself is not a complete reply to a nonstatutory double patenting (NSDP) rejection. A complete reply requires that the terminal disclaimer be accompanied by a reply requesting reconsideration of the prior Office action. Even where the NSDP rejection is provisional the reply must be complete. See MPEP § 804, subsection I.B.1. For a reply to a non-final Office action, see 37 CFR 1.111(a). For a reply to final Office action, see 37 CFR 1.113(c). A request for reconsideration while not provided for in 37 CFR 1.113(c) may be filed after final for consideration. See MPEP §§ 706.07(e) and 714.13.
The USPTO Internet website contains terminal disclaimer forms which may be used. Please visit www.uspto.gov/patent/patents-forms. The actual filing date of the application in which the form is filed determines what form (e.g., PTO/SB/25, PTO/SB/26, PTO/AIA /25, or PTO/AIA /26) should be used. A web-based eTerminal Disclaimer may be filled out completely online using web-screens. An eTerminal Disclaimer that meets all requirements is auto-processed and approved immediately upon submission. For more information about eTerminal Disclaimers, refer to www.uspto.gov/patents/apply/applying-online/eterminal-disclaimer.
Claim 1-36 provisionally rejected on the ground of nonstatutory double patenting as being unpatentable over claim 1, 10, 19 and 28 of copending Application No. 17/955,143 in view of Harris and Nickolls (US 7788468 B1)
17/955,153
17/955,143
Claim 1. One or more processors comprising: circuitry to, in response to a call to an application programming interface (API), cause one or more threads within a cluster comprising a plurality of groups of one or more blocks of threads to: execute a barrier instruction; and stop at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 1. One or more processors, comprising: circuitry to, in response to a call to s indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads.
Claim 2. The one or more processors of claim 1,wherein the one or more blocks of threads are of a software kernel.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 3. The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 4. The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 5 . The one or more processors of claim 1, wherein the barrier instruction is comprised within each thread that is comprised within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 6 . . The one or more processors of claim 1, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by one or more the threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 7 . The one or more processors of the circuitry to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 8 . The one or more processors of claim 1, the circuitry to, in response to the call to the API provide an indication that the barrier instruction was not executed by any of the one or more threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 9 . The one or more processors of claim 1, wherein the API is called by a thread within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 10 . A computer-implemented method comprising: receiving an application programming interface (API)call; executing a barrier instruction by one or more threads within a cluster, the cluster comprising a plurality of groups of one or more blocks of threads; and stopping the one or more threads within the cluster, at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 10. A computer-implemented method comprising: in response to a call to an application programming interface (API),indicating whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads.
Claim 11 . The computer-implemented method of claim 10, where in the one or more blocks of threads are of a software kernel .
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 12 . The computer-implemented method of claim 10, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 13 . The computer-implemented method of claim 10, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 14 . The computer-implemented method of claim 10, wherein the barrier instruction is comprised within each thread that is comprised within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 15 . The computer-implemented method of claim 10, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or more threads within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 16 . The computer-implemented method of claim 10, further comprising: indicating, in response to the API call, that the barrier instruction was executed by the one or more threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 17 . The computer-implemented method of claim 10, further comprising: indicating, in response to the API call, that the barrier instruction was not executed by any threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 18 . The computer-implemented method of claim 10 further comprising: executing the API call by a calling thread of the one or more threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 19 . A computer system comprising: one or more processors; and memory storing executable instructions that, when executed by the one or more processors, are to: in response to a call to an application programming interface (API) cause one or more threads within a cluster comprising a plurality of groups of one or more blocks of threads to: execute a barrier instruction; and stop at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 19 . A computer system comprising: one or more processors and at least one memory storing executable instructions that, when executed by the one or more processors, are to in response to a call to an application programming interface (API). indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by at least one thread of each of the two or more blocks of threads.
Claim 20 . The computer system of claim 19, wherein the one or more blocks of threads are of a software kernel.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 21 The computer system of claim 19, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 22 . The computer system of claim 19, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 23 . The computer system of claim 19, wherein the barrier instruction is comprised within in each thread that is comprised within the cluster
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 24 . The computer system of claim 19, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or nore threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 25 . The computer system of claim 19, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 26 . The computer system of claim 19, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was not executed by any of the one or more threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 27 . The computer system of claim 19, wherein the API is
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 28 . A non-transitory machine-readable medium having stored thereon a set of instructions, which if performed by one or more processors, are to: in response to a call to
Claim 28 . A non-transitory machine-readable medium having stored thereon a set of instructions, that when performed by one or more processors, are to in response to a call to an application programming interface (API)s indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads.
Claim 29 . The non-transitory machine-readable medium of claim 28, wherein the one or more blocks of threads are of a software kernel.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 30 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster group is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 31 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 32 . The non-transitory machine-readable medium of claim 28, wherein the barrier instruction is comprised within in each thread that is comprised within the cluster
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 33 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on performance of the barrier instruction by the threads within the cluster machine-readable medium of claim 28, wherein the group of blocks of threads is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or more threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 34 . The non-transitory machine-readable medium of claim 28, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 35 . The non-transitory machine-readable medium of claim 28,the one or more processors to, in response to the call to the API
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 36 . The non-transitory machine-readable medium of claim 28, wherein the API is
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
This is a provisional nonstatutory double patenting rejection.
Claims 1-36 provisionally rejected on the ground of nonstatutory double patenting as being unpatentable over claims 1, 10, 19 and 28 of copending Application No. 17/955,163 in view of Harris and Nickolls (US 7788468 B1)
This is a provisional nonstatutory double patenting rejection.
17/955,153
17/955,163
Claim 1. One or more processors comprising: circuitry to, in response to a call to an application programming interface (API), cause one or more threads within a cluster comprising a plurality of groups of one or more blocks of threads to: execute a barrier instruction; and stop at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 1. One or more processors, comprising: circuitry to, in response to a call an application programming interface (API)s indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads
Claim 2. The one or more processors of claim 1,wherein the one or more blocks of threads are of a software kernel.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 3. The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 4. The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 5 . The one or more processors of claim 1, wherein the barrier instruction is comprised within each thread that is comprised within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 6 . . The one or more processors of claim 1, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by one or more the threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 7 . The one or more processors of the circuitry to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 8 . The one or more processors of claim 1, the circuitry to, in response to the call to the API provide an indication that the barrier instruction was not executed by any of the one or more threads within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 9 . The one or more processors of claim 1, wherein the API is called by a thread within the cluster.
Claim 1. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 10 . A computer-implemented method comprising: receiving an application programming interface (API)call; executing a barrier instruction by one or more threads within a cluster, the cluster comprising a plurality of groups of one or more blocks of threads; and stopping the one or more threads within the cluster, at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 10. A computer-implemented method comprising: in response to a call to an application programming interface (API),to indicating whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads.
Claim 11 . The computer-implemented method of claim 10, where in the one or more blocks of threads are of a software kernel .
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 12 . The computer-implemented method of claim 10, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 13 . The computer-implemented method of claim 10, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 14 . The computer-implemented method of claim 10, wherein the barrier instruction is comprised within each thread that is comprised within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 15 . The computer-implemented method of claim 10, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or more threads within the cluster
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 16 . The computer-implemented method of claim 10, further comprising: indicating, in response to the API call, that the barrier instruction was executed by the one or more threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 17 . The computer-implemented method of claim 10, further comprising: indicating, in response to the API call, that the barrier instruction was not executed by any threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 18 . The computer-implemented method of claim 10 further comprising: executing the API call by a calling thread of the one or more threads within the cluster.
Claim 10. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 19 . A computer system comprising: one or more processors; and memory storing executable instructions that, when executed by the one or more processors, are to: in response to a call to an application programming interface (API) cause one or more threads within a cluster comprising a plurality of groups of one or more blocks of threads to: execute a barrier instruction; and stop at least until all threads within each of the plurality of groups of the cluster have performed corresponding barrier instructions
Claim 19 . A computer system comprising: one or more processors and at least one memory storing executable instructions that, when executed by the one or more processors, are to in response to a call to an application programming interface (API). indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by at least one thread of each of the two or more blocks of threads
Claim 20 . The computer system of claim 19, wherein the one or more blocks of threads are of a software kernel.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 21 The computer system of claim 19, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 22 . The computer system of claim 19, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 23 . The computer system of claim 19, wherein the barrier instruction is comprised within in each thread that is comprised within the cluster
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 24 . The computer system of claim 19, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or nore threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 25 . The computer system of claim 19, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 26 . The computer system of claim 19, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was not executed by any of the one or more threads within the cluster.
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 27 . The computer system of claim 19, wherein the API is
Claim 19. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 28 . A non-transitory machine-readable medium having stored thereon a set of instructions, which if performed by one or more processors, are to: in response to a call to
Claim 28 . A non-transitory machine-readable medium having stored thereon a set of instructions, that when performed by one or more processors, are to in response to a call to an application programming interface (API)s indicate whether one or more threads within two or more blocks of threads have performed a barrier instruction to prevent performance of one or more other instructions until performance of barrier instructions by threads of each of the two or more blocks of threads.
Claim 29 . The non-transitory machine-readable medium of claim 28, wherein the one or more blocks of threads are of a software kernel.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 30 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster group is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 31 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 32 . The non-transitory machine-readable medium of claim 28, wherein the barrier instruction is comprised within in each thread that is comprised within the cluster
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 33 . The non-transitory machine-readable medium of claim 28, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on performance of the barrier instruction by the threads within the cluster machine-readable medium of claim 28, wherein the group of blocks of threads is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose performance is not dependent on execution of the barrier instruction by the one or more threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 34 . The non-transitory machine-readable medium of claim 28, the one or more processors to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster.
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 35 . The non-transitory machine-readable medium of claim 28,the one or more processors to, in response to the call to the APIexecuted by any of the one or more threads within the cluster
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim 36 . The non-transitory machine-readable medium of claim 28, wherein the API is
Claim 28. In view of Nickolls. it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention to combine the teachings of Nickolls to improve the synchronizing of threads when executing multiple tasks, see art rejection below.
Claim Rejections - 35 USC § 103
In the event the determination of the status of the application as subject to AIA 35 U.S.C. 102 and 103 (or as subject to pre-AIA 35 U.S.C. 102 and 103) is incorrect, any correction of the statutory basis (i.e., changing from AIA to pre-AIA ) for the rejection will not be considered a new ground of rejection if the prior art relied upon, and the rationale supporting the rejection, would be the same under either status.
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.
Claims 1-36 are rejected under 35 U.S.C. 103 as being unpatentable over Harris “Cooperative Groups: Flexible CUDA Thread Programming” in view of Nickolls (US 7788468 B1).
Regarding claim 1, Harris teaches:
One or more processors comprising: circuitry to, in response to a call to an application programming interface (API), one or more threads within a cluster comprising a plurality of groups of one or more blocks of threads to: execute a barrier instruction;. (Page 2, The Cooperative Groups programming model describes synchronization patterns both within and across CUDA thread blocks. It provides CUDA device code APIs for defining, partitioning, and synchronizing groups of threads. It also provides host-side APIs to launch grids whose threads are all guaranteed to be executing concurrently to enable synchronization across thread blocks. These primitives enable new patterns of cooperative parallelism within CUDA, including producer-consumer parallelism and global synchronization across the entire thread grid or even multiple GPUs. Page 3, At its simplest, Cooperative Groups is an API for defining and synchronizing groups of threads in a CUDA program. Much of the Cooperative Groups (in fact everything in this post) works on any CUDA-capable GPU compatible with CUDA 9. Specifically, that means Kepler and later GPUs (Compute Capability 3.0+). Page 5, You can synchronize a group by calling its collective sync() method, or by calling the cooperative_groups::sync() function. These perform barrier synchronization among all threads in the group (Figure 2). See also page 6, “Thread Blocks”, page 8, “Partitioning Groups” and Figure 3 )
Harris does not appear to explicitly teach to stop at least until all threads within the cluster have performed a barrier instruction.
However, Nickolls teaches: col 2, line 29-35. Different threads of the CTA are advantageously synchronized at appropriate points during CTA execution using a barrier synchronization technique in which barrier instructions in the CTA program are detected and used to suspend execution of some threads until a specified number of other threads also reaches the barrier point. Col 22, line 53-67 Barrier synchronization logic 1130 includes a set of B counters 1134, a set of B target registers 1135, a comparison circuit 1136, and a set of wait/go registers 1138. Counters 1134 track the number of threads that have arrived at each of the B barrier points. Target registers 1105 store a target value associated with each of the B barrier points; in each case, the target value corresponds to the number of SIMD groups that are expected to synchronize at that barrier point. As in barrier synchronization logic 1100, the target value can be supplied as an immediate operand with the barrier instruction and is loaded into the appropriate target register 1135 by barrier detection circuit 1132 when the first barrier instruction pertaining to a particular barrier identifier BarID is received. Each target value remains stored in target register 1135 until its barrier is reset. See also col 23, line 33 – col 24, line 8.
Accordingly, it would have been obvious to a person having ordinary skill in the art before the effective filing date of the claimed invention, having the teachings of Harris and Nickolls before them, to include Nickolls’s barrier logic to stop a threat in Harris’s synchronization API for thread groups. One would have been motivated to apply the know technique of barrier mechanisms to enhance performance and ensure synchronization across multiple groups of threads an expected and predictable result in thread coordination.
Regarding claim 2, Harris teaches:
The one or more processors of claim 1, wherein the one or more blocks of threads are of a software kernel. (Page 6, Thread Blocks If you have programmed with CUDA before, you are familiar with thread blocks, the fundamental unit of parallelism in a CUDA program. Cooperative Groups introduces a new datatype, thread_block, to explicitly represent this concept within the kernel. An instance of thread_block is a handle to the group of threads in a CUDA thread block that you initialize as follows. thread_block block = this_thread_block();)
Regarding claim 3, Harris teaches:
The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads within the cluster. (Page 8, Partitioning Groups Cooperative Groups provides you the flexibility to create new groups by partitioning existing groups. This enables cooperation and synchronization at a finer granularity. The cg::tiled_partition() function partitions a thread block into multiple “tiles”. Here’s an example that partitions each whole thread block into tiles of 32 threads. thread_group tile32 = cg::tiled_partition(this_thread_block(), 32). See also figure 3.)
Regarding claim 4, Harris teaches:
The one or more processors of claim 1, wherein a group of the cluster is a partition of a partitioning of one or more blocks of threads of a grid of blocks of threads, the partitioning comprising multiple partitions. (Page 8, Partitioning Groups Cooperative Groups provides you the flexibility to create new groups by partitioning existing groups. This enables cooperation and synchronization at a finer granularity. The cg::tiled_partition() function partitions a thread block into multiple “tiles”. Here’s an example that partitions each whole thread block into tiles of 32 threads. thread_group tile32 = cg::tiled_partition(this_thread_block(), 32). See also figure 3 and page 6 Thread Blocks.)
Regarding claim 5, Harris teaches:
The one or more processors of claim 1, wherein the barrier instruction is comprised within each thread that is comprised within the cluster. (Page 4. Thread groups provide the ability to perform collective operations among all threads in a group. Collective operations, or simply collectives, are operations that need to synchronize or otherwise communicate amongst a specified set of threads. Because of the need for synchronization, every thread that is identified as participating in a collective must make a matching call to that collective operation. The simplest collective is a barrier, which transfers no data and merely synchronizes the threads in the group. Synchronization is supported by all thread groups)
Regarding claim 6, Nickolls teaches:
The one or more processors of claim 1, wherein a group of the cluster is in a grid of blocks of threads, and wherein the grid of blocks of threads comprises at least one block of one or more threads whose execution is not dependent on performance of the one or more barrier instruction by the threads within the cluster . (col 9, line 26-31. While all threads within a CTA are executed concurrently, there is no requirement that different CTAs are executed concurrently, and the hardware need not support sharing of data between threads in different CTAs. Thus, CTAs can be executed using any processing hardware with one or more processing engines. See also col 24, line 27-62)
Same motivation as claim 1.
Regarding claim 7, Nickolls teaches:
The one or more processors of the circuitry to, in response to the call to the API provide an indication that the barrier instruction was executed by the one or more threads within the cluster. (col 21, lines 33-50. Barrier synchronization logic 1100 includes a counter 1104, a target register 1105, a comparison circuit 1106, and wait/go registers 1108. Counter 1104 tracks the number of threads that have arrived at a particular barrier point. Target register 1105 stores a target value, which corresponds to the number of SIMD groups that are expected to synchronize at the barrier point. In one embodiment, the target value is supplied as an immediate operand with the barrier instruction and is loaded into target register 1105 by barrier detection circuit 1112 when the first barrier instruction is received. Once loaded, the target value advantageously remains stored in target register 1105 until target register 1105 is reset. Comparison circuit 1106 determines whether the number of arriving threads counted by counter 1104 has reached the target value stored in target register 1105. If the target value has been reached, comparison circuit 1106 issues a reset signal to counter 1104, target register 1105 and wait/go registers 1108.)
Same motivation as claim 1.
Regarding claim 8, Nickolls teaches:
The one or more processors of claim 1, the circuitry to, in response to the call to the API provide an indication that the barrier instruction was not executed by any of the one or more threads within the cluster. (col 22, line 10-24. Comparison circuit 1106 compares the current value in counter 1104 to the target value stored in register 1105. If the current value matches the target value, then the threads are properly synchronized and execution of any waiting threads can resume. Accordingly, comparison circuit 1106 generates a reset signal. The reset signal resets counter 1104 to zero, resets target register 1105 to an "unloaded" state (so that a new target value can be read in when the next barrier instruction is encountered), and resets wait/go registers 1108 such that the bits corresponding to all SIMD groups are in the go state. Selection unit 1110 thereafter resumes selecting instructions for the SIMD groups that were formerly waiting at the barrier point, allowing execution of those groups to proceed beyond the barrier point.)
Same motivation as claim 1.
Regarding claim 9, Harris teaches:
The one or more processors of claim 1, wherein the API is called by a thread within the cluster. (Page 5, You can synchronize a group by calling its collective sync() method, or by calling the cooperative_groups::sync() function. These perform barrier synchronization among all threads in the group (Figure 2). g.sync(); // synchronize group g cg::synchronize(g); // an equivalent way to synchronize g Here’s a simple example of a parallel reduction device function written using Cooperative Groups. When the threads of a group call it, they cooperatively compute the sum of the values passed by each thread in the group (through the val argument).)
With respect to claims 10-18: Claims 10-18 are directed to a method corresponding to the active functions implemented by the processor disclosed in claims 1-9, respectively; please see the rejections directed to claims 1-9 above which also cover the limitations recited in claims 10-18.
With respect to claims 19-27: Claims 19-27 are directed to a computer system comprising one or more processors and memory storing executable instructions to implement active functions corresponding to the active functions implemented by the processor disclosed in claims 1-9, respectively; please see the rejections directed to claims 1-9 above which also cover the limitations recited in claims 19-27.
With respect to claims 28-36: Claims 28-36 are directed to a machine-readable medium having stored thereon a set of instructions to implement active functions corresponding to the active functions implemented by the processor disclosed in claims 1-9, respectively; please see the rejections directed to claims 1-9 above which also cover the limitations recited in claims 28-36.
Conclusion
Any inquiry concerning this communication or earlier communications from the examiner should be directed to CARLOS A ESPANA whose telephone number is (703)756-1069. The examiner can normally be reached Monday - Friday 8 a.m - 5 p.m EST.
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, LEWIS BULLOCK JR can be reached at (571)272-3759. 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.
/C.A.E./Examiner, Art Unit 2199
/LEWIS A BULLOCK JR/Supervisory Patent Examiner, Art Unit 2199
Prosecution Insights