Few-shot training of neural networks

Abstract

This invention discloses a technique for training a neural network with a small sample size. The neural network is trained to recognize one or more features of an image. The neural network is trained using a small number of original images, from which multiple additional images are derived. The additional images are generated by rotating and decoding the image embeddings in a latent space generated by an autoencoder. The images generated by rotating and decoding exhibit changes in features proportional to the amount of rotation.

Description

Background Art

[0001] Few-shot learning is a training technique for neural networks that can be useful when it is difficult to obtain a sufficient number of training samples. Few-shot learning is applied to classification problems, where a neural network is trained to classify objects. However, few-shot learning has not generally been applied to regression problems. One such problem involves estimating the direction of a person's gaze from an input image that includes a representation of a person. BRIEF DESCRIPTION OF THE DRAWINGS

[0002] Figure 1 shows few-shot training for a regression problem according to one embodiment;

[0003] Figure 2 shows an example of the encoder portion of an autoencoder according to one embodiment;

[0004] Figure 3 shows an example of the rotation and decoding portion of an autoencoder according to one embodiment;

[0005] Figure 4 shows an example process for training an autoencoder according to one embodiment;

[0006] Figure 5 An example process for few-shot training of a neural network to solve a regression problem is shown according to one embodiment;

[0007] Figure 6 shows an example of a parallel processing unit ("PPU") according to one embodiment;

[0008] Figure 7 shows an example of a general processing cluster ("GPC") according to one embodiment;

[0009] Figure 8 An example of a memory partition unit according to one embodiment is shown;

[0010] Figure 9 shows an example of a streaming multiprocessor according to one embodiment; and

[0011] Figure 10 A computer system is shown in which various examples may be implemented according to one embodiment. DETAILED DESCRIPTION

[0012] Described herein are embodiments of a processor, system, method, and computer program product for training one or more neural networks to recognize features in an image. In one example embodiment, the processor is configured to recognize one or more features within an image based on a neural network trained using images generated from encoded points and rotated points in one or more equivariant latent spaces. The recognized features may include numerical estimates computed by the neural network as a solution to a regression problem.

[0013] In one embodiment, a neural network is trained using few-shot learning to solve a regression problem. As used herein, few-shot learning involves training a neural network to solve a regression problem using a relatively small training set. As an example of a regression problem, consider training a neural network to determine where a person is looking. This can be called view direction and can be considered a composite of gaze direction, which is sometimes also referred to as gaze angle, head rotation, and other possible factors. To estimate view direction, a neural network can be trained to estimate a person's gaze angle, head rotation angle, and so on based on images of the person. These estimates can also be called predictions. Regression problems may typically require a large training set, including examples covering the entire range of possible values. However, collecting such training sets can be difficult, especially when the accuracy of the solution requires personalized training. For example, with respect to the eye gaze problem, various differences in physical anatomy make a universal solution difficult. To improve accuracy, training can be personalized, but collecting such training sets from end users can be problematic. However, in the embodiments described herein, the few-shot training process can be applied to solve these and other regression problems. For example, embodiments of the techniques described herein may be applicable to medical diagnostics, audio, or other applications involving the training of neural networks about characteristics or traits of a particular individual.

[0014] In one embodiment, a neural network is trained to solve a regression problem based on training images generated by an autoencoder that is trained to embed images into an equivariant latent space. The generated images are sometimes referred to as hallucinated training images because they are artificially generated. Here, the training images can be based on a relatively small number of original images or real images. In some cases and embodiments, as few as one real image is used. Hallucinated images are generated by training an autoencoder to generate training images in which the attribute of interest (e.g., gaze direction) varies continuously across a range of values.

[0015] In one embodiment, a neural network is trained to solve a regression problem based on a code generated by an autoencoder that is trained to embed images into an equivariant latent space. This original code can be rotated to produce variants of the original code. These variants can then be used to train a neural network to solve the regression problem.

[0016] In one embodiment, an autoencoder is taught to embed features of an image into one or more equivariant latent spaces. As used herein, an equivariant latent space has the property that if the embedding of an image in the equivariant latent space is rotated by an amount r, decoding the rotated embedding will produce an image in which the attribute of interest changes by an amount proportional to r.

[0017] Figure 1 Figure 2 shows a small sample training for a regression problem according to one embodiment. Figure 1 In example 100 , the autoencoder 102 generates variant images or variant codes for use in the training process 108 .

[0018] In one embodiment, the autoencoder 102 is pre-trained to convert the input image x a is embedded into a latent space equivariant around the attribute of interest. The autoencoder 102 may include an encoder 104 and a decoder 106. Various aspects of training autoencoders are discussed in more detail herein, including Figure 4 .

[0019] In one embodiment, the input image x a is one of the relatively small number of images collected during the few-shot training process. Collect the input image x a , so that we know the image x a Once collected, the encoder 104 converts the image x into a Embedded into equivariant space. This embedding is then rotated and decoded multiple times to generate hallucinated images The hallucinated image shows a wide range of values ​​for the attribute of interest. In an embodiment, the process of rotation and decoding causes the image The values ​​of the attribute of interest are displayed across a desired range and with sufficient granularity to be used in training a neural network to solve a regression problem. In some cases and embodiments, each iteration of rotation and decoding includes an amount of rotation that is selected to produce sufficiently fine-grained differences in the hallucinated image to effectively train the network. Note that although Example 100 describes a single input image x a, but embodiments may employ a greater number of images. For example, in one embodiment, a 360° range of the attribute of interest is obtained by collecting one "real" image for each 90° quadrant. This example is intended to illustrate potential embodiments, not to limit, and therefore should not be interpreted as limiting the scope of the present disclosure to only those embodiments including the specific examples provided.

[0020] In one embodiment, the hallucinated image Each exhibits a value for the property of interest that is at least approximately proportional to the rotation in the equivariant latent space. This allows the image to be Each image in is labeled with a corresponding attribute value for use during draining. In an embodiment, the labeling of the image is updated and saved for each rotation and retained for subsequent neural network training.

[0021] In one embodiment, the hallucinated image The training process 108 is used to train the neural network 110 to solve the regression problem. For example, the neural network 110 can be trained to generate an output y, which includes the input image x i Neural network 110 may include any of a variety of neural network architectures and may employ any of a variety of training techniques, as appropriate given the selected network architecture and the available set of generated images.

[0022] In one embodiment, instead of generating hallucinated images Or in addition to generating hallucinatory images In addition, the neural network 110 is trained using codes generated by rotating the embedding in the equivariant latent space. For example, one can train the neural network 110 based on the rotation of x a The neural network 110 is trained to generate an output y by embedding one or more attributes of into a latent space and rotating the code generated by the embedding, which includes the input image x i The estimated value of the property being displayed.

[0023] Figure 2 An example of an encoder portion of an autoencoder according to one embodiment is shown. In particular, example 200 shows an autoencoder for the problem of estimating the direction a person is looking. It will be understood that example 200 is provided to illustrate various aspects of potential embodiments of the present disclosure and, therefore, should not be interpreted in a manner that limits the scope of the present disclosure to only those embodiments including the specific examples provided.

[0024] In one embodiment, one or more autoencoders are trained to generate an equivariant latent space of properties related to the eye gaze problem. Note that while Figure 2 A single encoder 202 is described, but embodiments may use one or more encoders similar to the described encoder 202 .

[0025] In one embodiment, the image x a is provided to the encoder 202 or to one or more such encoders, which then convert the image x a Embedded into one or more equivariant latent spaces. In an embodiment, the encoder 202 also generates one or more non-equivariant latent spaces.

[0026] In one embodiment, encoder 202 generates appearance code 204, gaze direction code 206, and head rotation code 208. Appearance code 204 indicates that image x a Gaze direction code 206 indicates the direction of the subject's eyes, although in some embodiments, each eye can be coded separately. Head rotation code 208 indicates the direction of the subject's head.

[0027] In one embodiment, the encoder 202 embeds the gaze direction code 206 and the head rotation code 208 into one or more equivariant latent spaces. For example, in the depicted embodiment, the encoder 202 maps the gaze direction code 206 to a point in the latent space such that if the point is rotated by a certain amount r, then the image generated by decoding the rotated point exhibits a change in the attribute of interest (in this case, gaze direction) that is proportional to r. Similarly, the encoder 202 maps the head rotation code 208 to a point in the latent space such that if the point is rotated by r, then the image generated by decoding the rotated point exhibits a change in the head rotation that is proportional to r.

[0028] In one embodiment, the equivariant latent space is a three-dimensional latent space. To rotate a point in the equivariant latent space, the embedding in the space is multiplied by the three-dimensional rotation matrix. Because the latent space is equivariant, the resulting point preserves the properties of interest and can be decoded to generate an image in which the preserved properties have been rotated by a proportional amount.

[0029] Figure 3 An example of the rotation and decoding portions of a rotational autoencoder according to one embodiment is shown. Figure 3 Example 300 involves Figure 2The decoder portion of the autoencoder described herein and relates to the same problem of estimating the direction and direction a person is looking at. Thus, it will be appreciated that example 300 is also provided to illustrate various aspects of potential embodiments of the present disclosure and should therefore not be interpreted in a manner that limits the scope of the present disclosure to only those embodiments including the provided feature examples.

[0030] In one embodiment, the rotation corresponding to Figure 2 Depicted are gaze direction codes 306 and head rotation codes 308. The rotation process 312 may include a process of applying one or more three-dimensional rotation matrices to the codes 306, 308 output by the encoder.

[0031] In one embodiment, the codes 306, 308 are independently rotated. For example, the gaze direction code 306 can be rotated by 15°, and the head rotation code 308 can be rotated by 4°. In these and other examples, the rotations can be described as disentangled because, in these and other embodiments, the features represented by the codes can be independently rotated without significantly affecting the other features.

[0032] In one embodiment, the appearance code 304 is not rotated. For example, an image may be generated by decoding the rotated codes 306, 308 and the unrotated appearance code 304. one.

[0033] In one embodiment, the appearance code 304, the gaze direction code 306, and the head rotation code 308 are rotated and then decoded by the decoder 302. The rotation and decoding may then generate an output image One of the images Each image in can be expressed as the input image x a The images generated by this rotation and decoding can then be used to train a neural network to solve a regression problem based on the data generated by multiple rotations.

[0034] In one embodiment, instead of generating an image And use these images for training or in addition to generating images In addition to using these images for training, the rotation appearance code 304 or the head rotation code 308 is used to train a neural network to solve the regression problem.

[0035] Figure 4An example process for training an autoencoder according to one embodiment is shown. Although described as a series of operations, the described order should not be interpreted as limiting the scope of the present disclosure to only those embodiments that conform to the described order. For example, in some cases, the described operations may be performed in an order different from the described order, or may be performed in parallel, except where the order may be logically required. In some embodiments, some of the described operations may be omitted.

[0036] In one embodiment, the example process 400 is performed by a suitable computing device including, but not limited to, Figure 6 The parallel processing units shown in , or Figure 7 The general processing cluster shown in .

[0037] In one embodiment, at 402, an image for training an autoencoder is obtained and provided to the encoder. The image may be designated as x a , which has known values ​​for one or more attributes of interest. For example, an image x a may have known values ​​for gaze direction or head rotation. Then, the obtained image x a Encoded by the encoder part of the autoencoder.

[0038] In one embodiment, at 404, the output of the encoder is rotated by a determined amount r. This amount is related to the first image x a and the second image x b is proportional to the difference in the property of interest between . For example, if the object is at x a The gaze direction in the x is 45°, while in the b If the gaze direction in is 50°, then r can be equal to 5° or otherwise proportional to 5°.

[0039] In one embodiment, at 406, the rotated output is decoded and a resulting image is generated. image Based on one or more such rotated outputs, and also based on additional unrotated outputs. Figure 3 As shown, the decoder utilizes gaze direction codes and head rotation codes that are rotated and appearance codes that are not subject to rotation.

[0040] In one embodiment, at 408, the image reconstructed based on the rotation and decoding is compared with the second image x b The second image is known to show a change in the property of interest that is proportional to the amount of rotation in the latent space. For example, if x aThe properties of interest shown are p a , then in x b For the same attribute of interest p b =Kr, where r is the amount of rotation in the latent space and K is some constant.

[0041] In one embodiment, the output image with x b In one embodiment, at 410, the error signal generated by the comparison is used to train the autoencoder to generate variants of the input image in which the properties of interest are changed by rotation, but are otherwise preserved. For example, once trained, the output image The property of interest displayed should be approximately p b =Kr.

[0042] In one embodiment, x a and x b Both may be synthetic data generated by various possible techniques.In an embodiment, any number of synthetic images are generated.

[0043] In one embodiment, a subject-specific autoencoder is trained and used to generate labels to assign to the generated images. This can be used to solve x a and x b Potential uncertainty in gaze angle estimation.

[0044] In one embodiment, Figure 4 As described above, the autoencoder is trained before training on a small sample of specific users. Once trained, the autoencoder can be used to generate images to train a neural network to solve the regression problem.

[0045] In one embodiment, the training of the autoencoder is performed using a relatively large dataset that includes samples for multiple individuals. Thus, the autoencoder can generally be trained to generate embeddings in an equivariant latent space.

[0046] In one embodiment, the training of the autoencoder is based at least in part on a clustering term. For example, the loss function used to train the autoencoder can be defined as:

[0047] Loss = L recon +L clustering +L prop

[0048] In the above embodiment of the loss function, Lrecon is the reconstruction loss, and L prop Indicates the loss of the attribute of interest. L clustering The term involves clustering points in images of the same person.

[0049] Figure 5 An example process for small-sample training of a neural network to solve a regression problem according to one embodiment is shown. Although described as a series of operations, the described order should not be interpreted as limiting the scope of the present disclosure to only those embodiments that conform to the described order. For example, in some cases, the described operations may be performed in an order different from the described order, or may be performed in parallel, except where the order may be logically required. In some embodiments, some of the described operations may be omitted.

[0050] In one embodiment, the example process 500 is performed by a suitable computing device including, but not limited to, Figure 6 The parallel processing unit described in , or Figure 7 The general processing cluster shown in .

[0051] In one embodiment, an image of a subject is obtained at 502. The image is obtained in a manner that allows for the marking of attributes of interest displayed in the obtained image. For example, with respect to eye gaze, the image can be obtained while indicating that the subject is looking at a particular location.

[0052] In an embodiment, at 504, multiple additional images or codes are generated based on the obtained image. The images or codes can be generated by taking the embeddings of an autoencoder and rotating them. For example, the obtained image is encoded by the encoder portion of the autoencoder and then incrementally rotated. The rotated codes can be used to train the network without generating variant images, or can be decoded and used to generate additional images. In these and other embodiments, the labels of the generated images can be derived from the labels of the original image and the accumulated amount of rotations.

[0053] In one embodiment, images or codes are generated based on the acquired images to produce examples across the range of values ​​to be estimated. Taking gaze direction as an example, the original real image may show a gaze direction of 0°, and the hallucinated image or latent code may be generated in steps of 0.5°, resulting in hallucinated images showing gaze directions ranging from 0.5° to 89.5°. Other real images can be obtained and used to generate hallucinated images for other ranges. The use of this method can be determined experimentally or based on available images.

[0054] In one embodiment, a neural network is trained to solve a regression problem based at least in part on the codes derived from the rotations or images generated based on these codes, at 506. A variety of network architectures can be used, and the network can be trained using a variety of training techniques that depend on the generated images.

[0055] In one embodiment, at 508, a neural network is used to solve the regression problem.

[0056] In one embodiment, the regression problem involves the estimation of values ​​related to attributes such as gaze direction or head rotation.

[0057] In one embodiment, the output of the neural network is back-propagated to the autoencoder and used to refine the autoencoder's generation of images used to train the neural network.

[0058] In one embodiment, a processor including one or more arithmetic logic units (ALUs) is configured to identify one or more features within an image based at least in part on one or more neural networks trained using encoded orientations of the one or more features. The identified features may include numerical estimates of solutions to a regression problem computed by the one or more neural networks. As used herein, encoded orientations refer to embeddings in an equivariant latent space.

[0059] In one embodiment, training of one or more neural networks is performed using images generated based on the encoded orientation of one or more features, such as based on embeddings in an equivariant latent space. The embeddings are generated by a decoder that has been taught the equivariant latent space, such as by Figure 4 The process described above is accomplished by generating a plurality of encoded orientations, and generating an image based on the rotated embedding and decoding it. When the equivariant latent space is rotated by a certain amount, the corresponding attribute of interest is preserved in the rotation. Thus, the resulting image exhibits attribute values ​​that change proportionally to the amount of rotation in the latent space compared to before the rotation. Thus, the image generated based on the rotated embedding exhibits a change in at least one of the one or more features that is proportional to the amount of rotation of the embedding.

[0060] In an embodiment, an autoencoder is trained to encode an image, wherein a second image generated by rotating and decoding the encoding of the image exhibits a value of an attribute (e.g., the gaze direction of an object depicted in the image) that varies proportionally with the amount of rotation. A method for solving a regression problem using a neural network may include training an autoencoder in this manner. The method also includes using the autoencoder to generate multiple images based on rotations of the encoding of the input image, and then using these images to train one or more neural networks to estimate the value of the attribute.

[0061] In one embodiment, the neural network regresses differences from the reference image. For example, the neural network can be trained to estimate the difference between one or more attributes displayed in the reference image and those attributes in the input image.

[0062] In one embodiment, the processor is configured to estimate a property based at least in part on one or more neural networks trained using images generated from an embedding of a first image, wherein the property of the first image changes in proportion to the amount of rotation in a second image generated by rotating and decoding the embedding.

[0063] Figure 6 A parallel processing unit ("PPU") 600 is shown according to one embodiment. In one embodiment, PPU 600 is configured using machine-readable code that, if executed by the PPU, causes the PPU to perform some or all of the processes and techniques described throughout this disclosure. In one embodiment, PPU 600 is implemented on one or more integrated circuit devices and utilizes multithreading as a latency hiding technique for a multithreaded processor designed to process computer-readable instructions (also referred to as machine-readable instructions or simple instructions) in parallel on multiple threads. In one embodiment, a thread refers to an execution thread and is an instance of an instruction set configured to be executed by PPU 600. In one embodiment, PPU 600 is a graphics processing unit ("GPU") configured to implement a graphics rendering pipeline for processing three-dimensional ("3D") graphics data to generate two-dimensional ("2D") image data for display on a display device (e.g., a liquid crystal display (LCD) device). In one embodiment, PPU 600 is used to perform computations such as linear algebra operations and machine learning operations. Figure 6 Exemplary parallel processors are shown for illustrative purposes only and should be construed as non-limiting examples of processor architectures contemplated within the scope of the present disclosure, and any suitable processors may be employed in addition and / or in place thereof.

[0064] In one embodiment, one or more PPUs are configured to accelerate high-performance computing ("HPC"), data center, and machine learning applications. In one embodiment, PPU 600 is configured to accelerate deep learning systems and applications, including the following non-limiting examples: autonomous vehicle platforms, deep learning, high-precision speech, image, and text recognition systems, intelligent video analysis, molecular simulation, drug discovery, disease diagnosis, weather forecasting, big data analytics, astronomy, molecular dynamics simulation, financial modeling, robotics, factory automation, real-time language translation, online search optimization, and personalized user recommendations.

[0065] In one embodiment, the PPU 600 includes an input / output ("I / O") unit 606, a front-end unit 610, a scheduler unit 612, a work distribution unit 614, a hub 616, a crossbar switch ("Xbar") 620, one or more general processing clusters ("GPCs") 618, and one or more partition units 622. In one embodiment, the PPU 600 is connected to a host processor or other PPUs 600 via one or more high-speed GPU interconnects 608. In one embodiment, the PPU 600 is connected to a host processor or other peripheral devices via an interconnect 602. In one embodiment, the PPU 600 is connected to a local memory including one or more storage devices 604. In one embodiment, the local memory includes one or more dynamic random access memory ("DRAM") devices. In one embodiment, the one or more DRAM devices are configured and / or configurable as a high-bandwidth memory ("HBM") subsystem, and multiple DRAM dies are stacked within each device.

[0066] The high-speed GPU interconnect 608 may refer to a wire-based, multi-lane communication link used by a system to extend and include one or more PPUs 600 in conjunction with one or more CPUs, supporting cache coherency between the PPU 600 and the CPU, and CPU mastering. In one embodiment, data and / or commands are sent by the high-speed GPU interconnect 608 to / from other units of the PPU 600 (e.g., one or more copy engines, video encoders, video decoders, power management units, and other components not in the PPU 600) through the hub 616. Figure 6 other components explicitly shown in the ).

[0067] In one embodiment, the I / O unit 606 is configured to receive data from the host processor ( Figure 6606 sends and receives communications (e.g., commands, data) via the system bus 602. In one embodiment, the I / O unit 606 communicates with the host processor directly via the system bus 602 or through one or more intermediate devices (e.g., a memory bridge). In one embodiment, the I / O unit 606 can communicate with one or more other processors (e.g., one or more PPUs 600) via the system bus 602. In one embodiment, the I / O unit 606 implements a Peripheral Component Interconnect Express ("PCIe") interface for communicating via a PCIe bus. In one embodiment, the I / O unit 606 implements an interface for communicating with external devices.

[0068] In one embodiment, the I / O unit 606 decodes packets received via the system bus 602. In one embodiment, at least a portion of the packets represent commands configured to cause the PPU 600 to perform various operations. In one embodiment, the I / O unit 606 sends the decoded commands to various other units of the PPU 600 as specified by the commands. In one embodiment, the commands are sent to the front end unit 610 and / or to the hub 616 or other units of the PPU 600, such as one or more copy engines, video encoders, video decoders, power management units, etc. (in the example of FIG. 1 ). Figure 6 In one embodiment, I / O unit 606 is configured to route communications between and among the various logical units of PPU 600.

[0069] In one embodiment, a program executed by a host processor encodes a command stream in a buffer that provides a workload to PPU 600 for processing. In one embodiment, the workload includes instructions and data to be processed by those instructions. In one embodiment, the buffer is an area in memory that is accessible (e.g., read / write) to both the host processor and PPU 600—the host interface unit can be configured to access the buffer in system memory connected to system bus 602 via memory requests sent by I / O unit 606 over system bus 602. In one embodiment, the host processor writes a command stream into the buffer and then transmits a pointer to the beginning of the command stream to PPU 600, so that front end unit 610 receives pointers to one or more command streams, manages the one or more streams, reads commands from the streams, and forwards the commands to various units of PPU 600.

[0070] In one embodiment, the front end unit 610 is coupled to a scheduler unit 612, which configures the various GPCs 618 to process tasks defined by one or more streams. In one embodiment, the scheduler unit 612 is configured to track state information related to the various tasks managed by the scheduler unit 612, where the state information may indicate which GPC 618 the task is assigned to, whether the task is active or inactive, a priority associated with the task, etc. In one embodiment, the scheduler unit 612 manages the execution of multiple tasks on one or more GPCs 618.

[0071] In one embodiment, the scheduler unit 612 is coupled to a work distribution unit 614, which is configured to distribute tasks to be executed on the GPCs 618. In one embodiment, the work distribution unit 614 tracks a plurality of scheduled tasks received from the scheduler unit 612. The work distribution unit 614 also manages a pending task pool and an active task pool for each GPC 618. In one embodiment, the pending task pool includes a plurality of time slots (e.g., 32 time slots) containing tasks assigned to be processed by a particular GPC 618; the active task pool may include a plurality of time slots (e.g., 4 time slots) for tasks being actively processed by the GPC 618, such that when a GPC 618 completes execution of a task, the task is evicted from the active task pool of the GPC 618, and one of the other tasks from the pending task pool is selected and scheduled for execution on the GPC 618. In one embodiment, if an active task is idle on a GPC 618 , such as while waiting for a data dependency to be resolved, the active task is evicted from the GPC 618 and returned to the pending task pool while another task is selected from the pending task pool and scheduled for execution on the GPC 618 .

[0072] In one embodiment, work distribution unit 614 communicates with one or more GPCs 618 via XBar 620. In one embodiment, XBar 620 is an interconnect network that couples many units of PPU 600 to other units of PPU 600, and it can be configured to couple work distribution unit 614 to a specific GPC 618. Although not explicitly shown, one or more other units of PPU 600 can also be connected to XBar 620 through hub 616.

[0073] These tasks are managed by the scheduler unit 612 and assigned to GPCs 618 by the work distribution unit 614. The GPCs 618 are configured to process tasks and generate results. The results can be consumed by other tasks within the GPC 618, can be routed to another GPC 618 via the XBar 620, or can be stored in the memory 604. The results can be written to the memory 604 by the partition unit 622, which implements a memory interface for reading data from and writing data to the memory 604. The results can be transferred to another PPU 604 or the CPU via the high-speed GPU interconnect 608. In one embodiment, the PPU 600 includes a number U of partition units 622, which is equal to the number of independent and distinct memory devices 604 coupled to the PPU 600. Figure 6 The partition unit 622 is described in more detail.

[0074] In one embodiment, a host processor executes a driver kernel that implements an application programming interface ("API") that enables one or more applications executing on the host processor to schedule operations for execution on the PPU 600. In one embodiment, the PPU 600 executes multiple computing applications concurrently, and the PPU 600 provides isolation, quality of service ("QoS"), and independent address spaces for the multiple computing applications. In one embodiment, the application generates instructions (e.g., in the form of API calls) that cause the driver kernel to generate one or more tasks to be executed by the PPU 600, and the driver kernel outputs the tasks to one or more streams being processed by the PPU 600. In one embodiment, each task includes one or more related groups of threads, which may be referred to as warps. In one embodiment, a warp includes multiple related threads (e.g., 32 threads) that may be executed in parallel. In one embodiment, cooperating threads may refer to multiple threads that include instructions to perform tasks and exchange data through shared memory. According to one embodiment, in combination Figure 7 Describes threads and cooperative threads in more detail.

[0075] Figure 7 According to one embodiment, Figure 6In one embodiment, each GPC 700 includes multiple hardware units for processing tasks, and each GPC 700 includes a pipeline manager 702, a pre-raster operations unit ("PROP") 704, a raster engine 708, a work distribution crossbar ("WDX") 716, a memory management unit ("MMU") 718, one or more data processing clusters ("DPCs") 706, and any suitable combination of components. It will be understood that Figure 7 The GPC 700 may include instead of or in addition to Figure 7 Other hardware units of the unit shown.

[0076] In one embodiment, the operation of GPC 700 is controlled by pipeline manager 702. Pipeline manager 702 manages the configuration of one or more DPCs 706 to process tasks assigned to GPC 700. In one embodiment, pipeline manager 702 configures at least one of one or more DPCs 706 to implement at least a portion of a graphics rendering pipeline. In one embodiment, DPC 706 is configured to execute vertex shader programs on a programmable streaming multiprocessor ("SM") 714. Pipeline manager 702 is configured to route packets received from work assignments to appropriate logic units within GPC 700. In one embodiment, some packets may be routed to fixed-function hardware units in PROP 704 and / or raster engine 708, while other packets may be routed to DPC 706 for processing by primitive engine 712 or SM 714. In one embodiment, pipeline manager 702 configures at least one of one or more DPCs 706 to implement a neural network model and / or a compute pipeline.

[0077] In one embodiment, the PROP unit 704 is configured to route data generated by the raster engine 708 and the DPC 706 to a raster operations ("ROP") unit in the memory partitioning unit, as described in more detail above. In one embodiment, the PROP unit 704 is configured to perform optimizations for color blending, organize pixel data, perform address translation, etc. In one embodiment, the raster engine 708 includes a plurality of fixed-function hardware units configured to perform various raster operations, and the raster engine 708 includes a setup engine, a coarse raster engine, a culling engine, a clipping engine, a fine raster engine, a tile merging engine, and any suitable combination thereof. In one embodiment, the setup engine receives the transformed vertices and generates plane equations associated with the geometric primitives defined by these vertices; transmits the plane equations to the coarse raster engine to generate coverage information for the primitives (e.g., the x, y coverage mask for the tile); transmits the output of the coarse raster engine to the culling engine, where fragments associated with primitives that fail the z-test are culled, and to the clipping engine, where fragments outside the viewing frustum are clipped. In one embodiment, the fragments remaining after clipping and culling are passed to the fine raster engine to generate pixel fragment attributes based on the plane equations generated by the setup engine. In one embodiment, the output of the raster engine 708 includes fragments to be processed by any suitable entity (e.g., by a fragment shader implemented within the DPC 706).

[0078] In one embodiment, each DPC 706 included in a GPC 700 includes an M-Pipe Controller ("MPC") 710; a primitive engine 712; one or more SMs 714; and any suitable combination thereof. In one embodiment, the MPC 710 controls the operation of the DPC 706, routing packets received from the pipeline manager 702 to appropriate units within the DPC 706. In one embodiment, packets associated with vertices are routed to the primitive engine 712, which is configured to retrieve vertex attributes associated with the vertices from memory; conversely, packets associated with shader programs may be transferred to the SM 714.

[0079] In one embodiment, SM 714 includes a programmable streaming processor configured to process tasks represented by multiple threads. In one embodiment, SM 714 is multithreaded and configured to simultaneously execute multiple threads (e.g., 32 threads) from a particular thread group, implementing a SIMD (single instruction, multiple data) architecture, wherein each thread in a thread group (e.g., a warp) is configured to process a different data set based on the same instruction set. In one embodiment, all threads in a thread group execute the same instructions. In one embodiment, SM 714 implements a SIMT (single instruction, multiple thread) architecture, wherein each thread in a thread group is configured to process a different data set based on the same instruction set, but wherein individual threads in a thread group are allowed to diverge during execution. In one embodiment, a program counter, call stack, and execution state are maintained for each warp, thereby enabling concurrency between warps and serial execution within a warp when threads within the warp diverge. In another embodiment, a program counter, call stack, and execution state are maintained for each individual thread, thereby enabling equal concurrency between all threads within and between warps. In one embodiment, execution state is maintained for each individual thread, and threads executing the same instruction can be converged and executed in parallel for better efficiency. In one embodiment, SM 714 will be described in more detail below.

[0080] In one embodiment, MMU 718 provides an interface between GPC 700 and the memory partition unit, and MMU 718 provides virtual-to-physical address translation, memory protection, and arbitration of memory requests. In one embodiment, MMU 718 provides one or more translation lookaside buffers ("TLBs") for translating virtual addresses into physical addresses in memory.

[0081] Figure 8 1 shows a memory partition unit of a PPU according to one embodiment. In one embodiment, the memory partition unit 800 includes a raster operation ("ROP") unit 802; a level 2 ("L2") cache 804; a memory interface 806, and any suitable combination thereof. The memory interface 806 is coupled to a memory. The memory interface 806 can implement a 32-, 64-, 128-, 1024-bit data bus, etc., for high-speed data transfer. In one embodiment, the PPU includes U memory interfaces 806, one memory interface 806 for each pair of partition units 800, where each pair of partition units 800 is connected to a corresponding memory device. For example, the PPU can be connected to up to Y memory devices, such as a high-bandwidth memory stack or graphics double data rate version 5 synchronous dynamic random access memory ("GDDR5 SDRAM").

[0082] In one embodiment, memory interface 806 implements an HBM2 memory interface, and Y is equal to half of U. In one embodiment, the HBM2 memory stack is located on the same physical package as the PPU, saving significant power and area compared to traditional GDDR5 SDRAM systems. In one embodiment, each HBM2 stack includes four memory dies, and Y is equal to 4. The HBM2 stack includes two 128-bit channels per die, for a total of 8 channels, and a data bus width of 1024 bits.

[0083] In one embodiment, the memory supports single-error correction, double-error detection ("SECDED") error correction code ("ECC") to protect data. ECC provides increased reliability for computing applications that are sensitive to data corruption. Reliability is particularly important in large-scale cluster computing environments where the PPU processes very large data sets and / or long-running applications.

[0084] In one embodiment, the PPU implements a multi-level memory hierarchy. In one embodiment, the memory partitioning unit 800 supports unified memory to provide a single unified virtual address space for CPU and PPU memory, thereby enabling data sharing between virtual memory systems. In one embodiment, the frequency of PPU accesses to memory located on other processors is tracked to ensure that memory pages are moved to the physical memory of the PPU that accesses the pages more frequently. In one embodiment, the high-speed GPU interconnect 608 supports address translation services that allow the PPU to directly access the CPU's page tables and provide full access to CPU memory by the PPU.

[0085] In one embodiment, the copy engine transfers data between multiple PPUs or between a PPU and a CPU. In one embodiment, the copy engine can generate a page fault for an address that is not mapped in the page table, and the memory partition unit 800 then services the page fault, maps the address into the page table, and then the copy engine performs the transfer. In one embodiment, fixed memory (i.e., non-pageable) is operated for multiple copy engines between multiple processors, thereby substantially reducing the available memory. In one embodiment, due to hardware page faults, addresses can be passed to the copy engine without having to consider whether the memory page is resident, and the copy process is transparent.

[0086] According to one embodiment, from Figure 6Data from the memory or other system memory is retrieved by the memory partition unit 800 and stored in the L2 cache 804, which is located on-chip and shared between the various GPCs. In one embodiment, each memory partition unit 800 includes at least a portion of the L2 cache 804 associated with the corresponding storage device. In one embodiment, lower-level caches are implemented in various units within the GPC. In one embodiment, each SM 714 may implement a level 1 ("L1") cache, where the L1 cache is private memory dedicated to a particular SM 840, and data is retrieved from the L2 cache 804 and stored in each L1 cache for processing in the functional units of the SM 714. In one embodiment, the L2 cache 804 is coupled to the memory interface 806 and the XBar 620.

[0087] In one embodiment, ROP unit 802 performs graphics raster operations related to pixel color, such as color compression and pixel blending. In one embodiment, ROP unit 802 performs depth testing in conjunction with raster engine 708, receiving the depth of sample locations associated with pixel fragments from the culling engine of raster engine 708. In one embodiment, the depth of the sample location associated with the fragment is tested against the corresponding depth in the depth buffer. In one embodiment, if the fragment passes the depth test for the sample location, ROP unit 802 updates the depth buffer and sends the result of the depth test to raster engine 708. It will be appreciated that the number of partition units 800 can differ from the number of GPCs, and therefore, in one embodiment, each ROP unit 802 can be coupled to each GPC. In one embodiment, ROP unit 802 tracks packets received from different GPCs and determines to which GPC to route the results generated by ROP unit 802 via Xbar.

[0088] Figure 9 According to one embodiment, Figure 7Streaming multiprocessors such as the SM 900 of the present invention are described herein. In one embodiment, the SM 900 includes: an instruction cache 902; one or more scheduler units 904; a register file 908; one or more processing cores 910; one or more specific functional units ("SFUs") 912; one or more load / store units ("LSUs") 914; an interconnect network 916; a shared memory / L1 cache 918; and any suitable combination thereof. In one embodiment, the work distribution unit distributes tasks for execution on the GPCs of the PPU, with each task being assigned to a specific DPC within the GPC, and if the task is associated with a shader program, the task is distributed to the SM 900. In one embodiment, the scheduler unit 904 receives tasks from the work distribution unit and manages the scheduling of instructions for one or more thread blocks assigned to the SM 900. In one embodiment, the scheduler unit 904 schedules the thread blocks for execution as warps of parallel threads, wherein each thread block is assigned at least one warp. In one embodiment, each warp executes a thread. In one embodiment, the scheduler unit 904 manages multiple different thread blocks, assigns warps to different thread blocks, and then distributes instructions from multiple different cooperation groups to various functional units (e.g., core 910, SFU 912, and LSU 914) during each clock cycle.

[0089] Cooperative Groups refers to a programming model for organizing groups of communicating threads. This programming model allows developers to express the granularity of communicating threads, enabling richer, more efficient decompositions of parallelism. In one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In one embodiment, applications of conventional programming models provide a single, simple construct for synchronizing cooperating threads: a barrier across all threads of a thread block (e.g., the syncthreads() function). However, programmers often want to define thread groups at a granularity smaller than the thread block and synchronize within the defined group to achieve higher performance, design flexibility, and software reuse in the form of collective group-wide functional interfaces. Cooperative Groups enables programmers to explicitly define thread groups at sub-block (i.e., down to a single thread) and multi-block granularity and perform collective operations, such as synchronization, on threads within the cooperative group. The programming model supports clean composition across software boundaries, so libraries and utility functions can safely synchronize within their local context without making assumptions about convergence. The Cooperative Groups primitive enables new cooperative parallelism patterns, including producer-consumer parallelism, opportunistic parallelism, and global synchronization across an entire grid of thread blocks.

[0090] In one embodiment, the dispatch unit 906 is configured to transfer instructions to one or more functional units, and the scheduler unit 904 includes two dispatch units 906, which enable dispatching two different instructions from the same warp during each clock cycle. In one embodiment, each scheduler unit 904 includes a single dispatch unit 906 or additional dispatch units 906.

[0091] In one embodiment, each SM 900 includes a register file 908 that provides register sets for the functional units of SM 900. In one embodiment, register file 908 is partitioned between each functional unit so that each functional unit is allocated a dedicated portion of register file 908. In one embodiment, register file 908 is partitioned between different warps being executed by SM 900, and register file 908 provides temporary storage for operands connected to the data paths of the functional units. In one embodiment, each SM 900 includes a plurality of L processing cores 910. In one embodiment, SM 900 includes a large number (e.g., 128 or more) of different processing cores 910. In one embodiment, each core 910 includes a fully pipelined, single-precision, double-precision, and / or mixed-precision processing unit, including a floating-point arithmetic logic unit and an integer arithmetic logic unit. In one embodiment, the floating-point arithmetic logic unit implements the IEEE 754-2008 standard for floating-point arithmetic. In one embodiment, core 910 includes 64 single-precision (32-bit) floating point cores, 64 integer cores, 32 double-precision (64-bit) floating point cores, and 8 tensor cores.

[0092] According to one embodiment, the tensor cores are configured to perform matrix operations. In one embodiment, one or more tensor cores are included in core 910. In one embodiment, the tensor cores are configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inference. In one embodiment, each tensor core operates on a 4×4 matrix and performs a matrix multiplication and accumulation operation D=A×B+C, where A, B, C, and D are 4×4 matrices.

[0093] In one embodiment, the matrix multiplication inputs A and B are 16-bit floating point matrices, and the accumulation matrices C and D are 16-bit floating point or 32-bit floating point matrices. In one embodiment, the tensor cores operate on 16-bit floating point input data with 32-bit floating point accumulation. In one embodiment, this 16-bit floating point multiplication requires 64 operations and produces a full-precision product, which is then added to other intermediate products using 32-bit floating point addition to perform a 4×4×4 matrix multiplication. In one embodiment, the tensor cores are used to perform larger two-dimensional or higher-dimensional matrix operations composed of these smaller elements. In one embodiment, an API such as the CUDA 9 C++ API exposes specialized matrix load, matrix multiplication and accumulation, and matrix store operations to efficiently use the tensor cores from CUDA-C++ programs. In one embodiment, at the CUDA level, the warp-level interface assumes a 16×16 size matrix across all 32 threads of the warp.

[0094] In one embodiment, each SM 900 includes M SFUs 912 that perform specific functions (e.g., attribute evaluation, reciprocal square root, etc.). In one embodiment, the SFUs 912 include tree traversal units configured to traverse a hierarchical tree data structure. In one embodiment, the SFUs 912 include texture units configured to perform texture map filtering operations. In one embodiment, the texture units are configured to load texture maps (e.g., 2D arrays of pixels) from memory and sample the texture maps to generate sampled texture values ​​for use in shader programs executed by the SM 900. In one embodiment, the texture maps are stored in shared memory / L1 cache. According to one embodiment, the texture units perform texture operations, such as filtering operations using mip-maps (e.g., texture maps with different levels of detail). In one embodiment, each SM 900 includes two texture units.

[0095] In one embodiment, each SM 900 includes N LSUs 914 that implement load and store operations between shared memory / L1 cache 918 and register file 908. In one embodiment, each SM 900 includes an interconnect network 916 that connects each functional unit to register file 908 and connects the LSUs 914 to register file 908 and shared memory / L1 cache 918. In one embodiment, the interconnect network 916 is a crossbar switch that can be configured to connect any functional unit to any register in register file 908 and to connect the LSUs 914 to storage locations in the register file and shared memory / L1 cache 918.

[0096] In one embodiment, shared memory / L1 cache 918 is an array of on-chip memory that allows data storage and communication between the SM 900 and the primitive engines, as well as between threads within the SM 900. In one embodiment, shared memory / L1 cache 918 includes 128KB of storage capacity and is located in the path from the SM 900 to the partition unit. In one embodiment, shared memory / L1 cache 918 is used to cache reads and writes. One or more of shared memory / L1 cache 918, L2 cache, and memory is a backing store.

[0097] In one embodiment, data cache and shared memory functionality are combined into a single memory block, providing improved performance for both types of memory accesses. In one embodiment, the capacity is used or available as a cache by programs that do not use the shared memory, for example, if the shared memory is configured to use half the capacity, texture and load / store operations can use the remaining capacity. According to one embodiment, the integration within the shared memory / L1 cache 918 enables the shared memory / L1 cache 918 to be used as a high throughput pipeline for streaming data, while providing high bandwidth and low latency access to frequently reused data. When configured for general parallel computing, a simpler configuration can be used compared to graphics processing. In one embodiment, the fixed function graphics processing unit is bypassed, creating a simpler programming model. In one embodiment, in a general parallel computing configuration, the work distribution unit directly allocates and distributes blocks of threads to the DPC. According to one embodiment, threads in a block execute the same program, use unique thread IDs in computations to ensure each thread generates unique results, use SM 900 to execute the program and perform computations, use shared memory / L1 cache 918 to communicate between threads, and use LSU 914 to read and write global memory through shared memory / L1 cache 918 and a memory partitioning unit. In one embodiment, when configured for general-purpose parallel computation, SM 900 writes commands that the scheduler unit can use to start new work on a DPC.

[0098] In one embodiment, the PPU is included in or coupled to a desktop computer, a laptop computer, a tablet computer, a server, a supercomputer, a smartphone (e.g., wireless, handheld device), a personal digital assistant ("PDA"), a digital camera, a vehicle, a head-mounted display, a handheld electronic device, etc. In one embodiment, the PPU is embodied on a single semiconductor substrate. In one embodiment, the PPU is included in a system-on-chip ("SoC") along with one or more other devices (e.g., an additional PPU, memory, a reduced instruction set computer ("RISC") CPU, a memory management unit ("MMU"), a digital-to-analog converter ("DAC"), etc.).

[0099] In one embodiment, the PPU may be included on a graphics card that includes one or more storage devices. The graphics card may be configured to interface with a PCIe slot on a desktop computer motherboard. In another embodiment, the PPU may be an integrated graphics processing unit ("iGPU") included in the motherboard's chipset.

[0100] Figure 10 A computer system 1000 is shown in which various architectures and / or functions may be implemented according to one embodiment. In one embodiment, the computer system 1000 is configured to implement the various processes and methods described throughout this disclosure.

[0101] In one embodiment, computer system 1000 includes at least one central processing unit 1002 connected to a communication bus 1010 implemented using any suitable protocol (e.g., PCI (Peripheral Component Interconnect), PCI-Express, AGP (Accelerated Graphics), HyperTransport) or any other bus or one or more point-to-point communication protocols. In one embodiment, computer system 1000 includes main memory 1004 and control logic (e.g., implemented as hardware, software, or a combination thereof), and data is stored in main memory 1004, which may take the form of random access memory ("RAM"). In one embodiment, network interface subsystem 1022 provides an interface to other computing devices and networks for receiving data from and sending data to other systems from computer system 1000.

[0102] In one embodiment, computer system 1000 includes an input device 1008, a parallel processing system 1012, and a display device 1006. Display device 1006 can be implemented using a conventional CRT (cathode ray tube), LCD (liquid crystal display), LED (light emitting diode), plasma display, or other suitable display technology. In one embodiment, user input is received from input device 1008 such as a keyboard, mouse, touchpad, microphone, etc. In one embodiment, each of the aforementioned modules can be located on a single semiconductor platform to form a processing system.

[0103] Throughout this specification, a single semiconductor platform may refer to a single, single semiconductor-based integrated circuit or chip. It should be noted that the term "single semiconductor platform" may also refer to a multi-chip module with increased connectivity that emulates on-chip operations and offers substantial improvements over traditional central processing unit ("CPU") and bus implementations. Of course, the various modules may also be placed separately or in various combinations within the semiconductor platform, depending on user needs.

[0104] In one embodiment, a computer program in the form of a machine-readable executable code or computer control logic algorithm is stored in the main memory 1004 and / or secondary storage. According to one embodiment, the computer program, if executed by one or more processors, enables the system 1000 to perform various functions. The memory 1004, memory and / or any other memory are possible examples of computer-readable media. Secondary storage can refer to any suitable storage device or system, such as a hard drive and / or removable storage drive, a floppy disk drive, a tape drive, an optical drive, a digital versatile disk ("DVD") drive, a recording device, a universal serial bus ("USB") flash memory.

[0105] In one embodiment, the various architectures and / or functions described in the previous figures may be implemented in the context of a central processing unit 1002, a parallel processing system 1012, an integrated circuit capable of simultaneously having at least a portion of the capabilities of the central processing unit 1002 and the parallel processing system 1012, a chipset (e.g., a group of integrated circuits designed to operate and sold as a unit performing related functions), and any suitable combination of integrated circuits.

[0106] In one embodiment, the architecture and / or functionality of the various previous figures are implemented in the context of a general purpose computer system, a circuit board system, a game console system dedicated for entertainment purposes, a dedicated system, etc. In one embodiment, the computer system 1000 may take the form of a desktop computer, a laptop computer, a tablet computer, a server, a supercomputer, a smartphone (e.g., a wireless handheld device), a personal digital assistant (“PDA”), a digital camera, a vehicle, a head mounted display, a handheld electronic device, a mobile telephone device, a television, a workstation, a game console, an embedded system, and / or any other type of logic.

[0107] In one embodiment, parallel processing system 1012 includes multiple PPUs 1014 and associated memory 1016. In one embodiment, the PPUs are connected to a host processor or other peripheral devices via an interconnect 1018 and a switch 1020 or multiplexer. In one embodiment, parallel processing system 1012 distributes computational tasks across parallelizable PPUs 1014—e.g., as part of a computational task distribution across multiple GPU thread blocks. In one embodiment, although memory is shared and accessed (e.g., for read and / or write access) between some or all of the PPUs 1014, such shared memory may incur a performance penalty relative to using local memory and registers resident on the PPUs. In one embodiment, the operation of the PPUs 1014 is synchronized using a command such as __syncthreads(), which requires all threads in a block (e.g., executing across multiple PPUs 1014) to reach a certain execution point in the code before continuing.

[0108] Accordingly, the specification and drawings are to be regarded in an illustrative rather than a restrictive sense.It will, however, be evident that various modifications and changes may be made thereto without departing from the broader spirit and scope of the invention as set forth in the claims.

[0109] Other variations are within the spirit of the present disclosure. Thus, while the disclosed technology is susceptible to various modifications and alternative constructions, certain illustrative embodiments are shown in the drawings and have been described above in detail. However, it should be understood that there is no intention to limit the invention to one or more specific forms disclosed, but on the contrary, it is intended to cover all modifications, alternative constructions, and equivalents that fall within the spirit and scope of the invention as defined by the appended claims.

[0110] Unless otherwise indicated herein or clearly contradicted by the context, in the context of describing the disclosed embodiments (particularly in the context of the appended claims), the use of the terms "one" and "an" and similar references should be interpreted as covering both singular and plural forms. Unless otherwise indicated, the terms "comprising," "having," "including," and "containing" should be interpreted as open-ended terms (i.e., meaning "including but not limited to"). The term "connected" refers to a physical connection when unmodified and should be understood as a place where part or all is contained in, connected to, or connected together, even if there is any intervention. Unless otherwise indicated herein, the enumeration of numerical ranges herein is only intended to be used as a shorthand method to refer to each individual value falling within the range, and each individual value is incorporated into the specification as if it were individually described herein. Unless otherwise indicated by the context or contradicted by the context, the use of the terms "group" (e.g., "group of items") or "subset" should be interpreted as a non-empty set including one or more members. In addition, unless otherwise indicated by the context or contradicted by the context, the term "subset" of a corresponding set does not necessarily represent a proper subset of the corresponding set, but a subset and a corresponding set may be equal.

[0111] Unless expressly indicated otherwise or clearly contradicted by context, conjunction language in phrases of the form "at least one of A, B, and C" or "at least one of A, B, and C" is otherwise understood in context as generally used to refer to an item, term, or the like that may be A or B or C, or any non-empty subset of the set of A and B and C. For example, in the illustrative example of a set having three members, the conjunction phrases "at least one of A, B, and C" and "at least one of A, B, and C" refer to any of the following sets: {A}, {B}, {C}, {A, B}, {A, C}, {B, C}, {A, B, C}. Thus, such conjunction language is generally not intended to imply that certain embodiments require that each of at least one of A, at least one of B, and at least one of C be present. Additionally, unless expressly indicated otherwise or contradicted by context, the term "plurality" denotes pluralism (e.g., "a plurality of items" means a plurality of items). The number of a plurality of items is at least two, but may be more when expressly indicated or by context. Further, unless stated otherwise or clear from context, the phrase "based on" means "based at least in part on" rather than "based solely on."

[0112] Unless otherwise indicated herein or clearly contradicted by the context, the operations of the processes described herein may be performed in any suitable order. In one embodiment, processes such as those described herein (or variations and / or combinations thereof) are performed under the control of one or more computer systems configured with executable instructions and are implemented as code (e.g., executable instructions, one or more computer programs, or one or more applications) that is executed together on one or more processors by hardware or a combination thereof. In one embodiment, for example, the code is stored on a computer-readable storage medium in the form of a computer program that includes multiple instructions that can be executed by one or more processors. In one embodiment, the computer-readable storage medium is a non-transitory computer-readable storage medium that does not include transient signals (e.g., propagated transient electrical or electromagnetic transmissions) but includes non-transitory data storage circuits (e.g., buffers, caches, and queues) in a transient signal transceiver. In one embodiment, code (e.g., executable code or source code) is stored on a set of one or more non-transitory computer-readable storage media (or other memory for storing executable instructions) having stored thereon executable instructions that, when executed by (i.e., as a result of being executed by) one or more processors of a computer system, cause the computer system to perform the operations described herein. In one embodiment, the set of non-transitory computer-readable storage media includes a plurality of non-transitory computer-readable storage media, and one or more individual non-transitory storage media in the plurality of non-transitory computer-readable storage media lack all of the code, while the plurality of non-transitory computer-readable storage media collectively store all of the code. In one embodiment, the executable instructions are executed such that different instructions are executed by different processors—for example, the non-transitory computer-readable storage media store instructions and a main CPU executes some of the instructions while a graphics processor unit executes other instructions. In one embodiment, different components of the computer system have separate processors, and the different processors execute different subsets of the instructions.

[0113] In one embodiment, accordingly, a computer system is configured to implement one or more services that individually or collectively perform the operations of the processes described herein, and such a computer system is configured with applicable hardware and / or software that enables the operations to be performed. Furthermore, a computer system implementing an embodiment of the present disclosure is a single device, and in another embodiment, it is a distributed computer system comprising multiple devices that operate in different ways, such that the distributed computer system performs the operations described herein and such that a single device cannot perform all of the operations.

[0114] Unless otherwise specified, the use of any and all examples or exemplary language (e.g., "such as") provided herein is intended merely to better illuminate embodiments of the invention and does not constitute a limitation on the scope of the invention. No language in the specification should be construed as indicating any non-claimed element as essential to the practice of the invention.

[0115] Embodiments of the present disclosure are described herein, including the best mode known to the inventor for implementing the present invention. By reading the foregoing description, variations of those embodiments will become apparent to those of ordinary skill in the art. The inventor expects that skilled persons will appropriately adopt such variations, and the inventor wishes to practice the embodiments of the present disclosure in a manner different from that specifically described herein. Therefore, the scope of the present disclosure includes all modifications and equivalents of the subject matter recited in the appended claims as permitted by applicable law. In addition, unless otherwise noted herein or clearly contradicted by the context, any combination of the above-mentioned elements in all their possible variations is encompassed by the scope of the present disclosure.

[0116] All references, including publications, patent applications, and patents, cited herein are hereby incorporated by reference to the same extent as if each reference were individually and specifically indicated to be incorporated by reference.

[0117] In the description and claims, the terms "coupled" and "connected," along with their derivatives, may be used. It should be understood that these terms may not be intended as synonyms for each other. Rather, in specific examples, "connected" or "coupled" may be used to indicate that two or more elements are in direct or indirect physical or electrical contact with each other. "Coupled" may also mean that two or more elements are not in direct contact with each other, but still cooperate or interact with each other.

[0118] Unless otherwise specifically stated, it is understood that throughout the specification, terms such as "process," "calculate," "compute," "determine," etc. refer to the actions and / or processes of a computer or computing system or similar electronic computing device that manipulates and / or transforms data represented as physical quantities (e.g., electronic quantities) in the computing system's registers and / or memories into other data similarly represented as physical quantities in the computing system's memories, registers, or other such information storage, transmission, or display devices.

[0119] In a similar manner, the term "processor" may refer to any device or part of a device that processes electronic data from registers and / or memory and converts the electronic data into other electronic data that can be stored in registers and / or memory. As non-limiting examples, a "processor" may be a central processing unit (CPU) or a graphics processing unit (GPU). A "computing platform" may include one or more processors. As used herein, a "software" process may include, for example, software and / or hardware entities that perform work over time, such as tasks, threads, and intelligent agents. Moreover, each process may refer to multiple processes that execute instructions continuously or intermittently, sequentially, or in parallel. The terms "system" and "method" are used interchangeably herein so long as the system can embody one or more methods and the method can be considered a system.

[0120] In this document, reference may be made to obtaining, acquiring, receiving analog or digital data or inputting it into a subsystem, computer system, or computer-implemented machine. The process of obtaining, acquiring, receiving, or inputting analog and digital data can be accomplished in a variety of ways, such as by receiving data as a parameter of a function call or a call to an application programming interface (API). In some implementations, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting data via a serial or parallel interface. In another implementation, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting data from a providing entity to an acquisition entity via a computer network. Reference may also be made to providing, outputting, transmitting, sending, or presenting analog or digital data. In various examples, the process of providing, outputting, transmitting, sending, or presenting analog or digital data can be accomplished by transmitting data as input or output parameters of a function call, an application programming interface, or an interprocess communication mechanism.

[0121] Although the above discussion sets forth example implementations of the described technology, other architectures may be used to implement the described functionality and are intended to be within the scope of this disclosure. In addition, although specific responsibilities are defined above for discussion purposes, the various functions and responsibilities may be allocated and divided in different ways depending on the circumstances.

[0122] Furthermore, although the subject matter has been described in language specific to structural features and / or methodological acts, it is to be understood that the subject matter defined in the appended claims is not necessarily limited to the specific features or acts described. Rather, the specific features and acts are disclosed as example forms of implementing the claims.

Claims

1. A processor, comprising: one or more arithmetic logic units (ALUs) that recognize one or more features within the image based at least in part on one or more neural networks trained using one or more encoded directions of the one or more features; wherein the one or more encoded directions are generated by an encoder that embeds the image into an equivariant latent space, and the one or more encoded directions are generated at least in part based on rotating the embedding of the image in the equivariant latent space by a certain amount. 2 . The processor of claim 1 , wherein identifying the one or more features comprises forming a numerical estimate of a value of an attribute in the image.

3. The processor of claim 1 , wherein the one or more encoded directions are used to generate images for training the one or more neural networks.

4. The processor of claim 1 , wherein the image generated based on the rotated embedding exhibits a change for at least one of the one or more features that is proportional to the amount of rotation of the embedding.

5. The processor of claim 1 , wherein the encoded direction is generated based at least in part on the encoder being trained to map inputs to points in the equivariant latent space.

6. The processor of claim 1 , wherein the neural network is trained to recognize one or more features in an image based at least in part on one or more images generated by rotating the encoded orientation.

7. A system comprising: one or more computers comprising one or more processors for training one or more neural networks to recognize one or more features in an image, the one or more neural networks trained using one or more encoded directions of the one or more features; wherein the one or more encoded directions are generated by an encoder that embeds the image into an equivariant latent space, and the one or more encoded directions are generated at least in part based on rotating the embedding of the image in the equivariant latent space by a certain amount.

8. The system of claim 7, wherein the identifying of the one or more features comprises forming a numerical estimate of a value of an attribute in the image.

9. The system of claim 7, wherein the one or more encoded directions are one or more embeddings in the equivariant latent space that can be rotated while preserving properties of interest.

10. The system of claim 7, wherein an image generated based on the rotated embedding exhibits a change for at least one of the one or more features that is proportional to the amount of rotation of the embedding.

11. The system of claim 7, wherein the encoder is trained at least in part based on comparing a first image predicted to exhibit a value of at least one of the one or more features to a second image known to exhibit the value, wherein the first image is generated by rotating and decoding an embedding in the equivariant latent space.

12. A machine-readable medium having stored thereon a set of instructions which, if executed by one or more processors, cause the one or more processors to at least: causing one or more neural networks to be trained to recognize one or more features in an image, the one or more neural networks being trained using the one or more encoded directions of the one or more features; in, The one or more encoded directions are generated by an encoder that embeds the image into an equivariant latent space, wherein the one or more encoded directions are generated based at least in part on rotating the embedding of the image in the equivariant latent space by a certain amount.

13. The machine-readable medium of claim 12, wherein identifying the one or more features comprises forming a numerical estimate of a value of an attribute in the image.

14. The machine-readable medium of claim 12, wherein an image is generated based at least in part on a rotation of the encoded orientation to train the one or more neural networks to change the orientation of the one or more features in proportion to the amount of rotation.

15. The machine-readable medium of claim 12, wherein the equivariant latent space preserves properties of interest upon rotation.

16. The machine-readable medium of claim 15 , wherein the encoder is trained at least in part based on a first image predicted to exhibit a value of at least one of the one or more features for a second image known to exhibit the value, wherein the first image is generated by rotating and decoding an embedding in the equivariant latent space.

17. A system comprising: Camera; as well as a processor for identifying a gaze direction of an object depicted in an image obtained by the camera, the gaze direction being identified based at least in part on one or more neural networks trained using one or more encoded directions of the gaze direction; wherein the one or more encoded directions are generated by an encoder that embeds the image into an equivariant latent space, wherein the one or more encoded directions are generated at least in part based on rotating the embedding of the image in the equivariant latent space by a certain amount.

18. The system of claim 17, wherein the encoded direction comprises a three-dimensional code indicating a gaze direction.

19. The system of claim 17, wherein the encoded direction comprises a three-dimensional code indicating head rotation.

20. The system of claim 17, wherein the one or more encoded directions are used to generate images for training the one or more neural networks.

21. The system of claim 20, wherein the encoded direction is generated by rotating at least one of a gaze direction code or a head rotation code.

22. The system of claim 21, wherein the one or more neural networks are trained at least in part based on images generated by decoding the encoded directions.

23. The system of claim 17, wherein the one or more neural networks are trained to solve for regression of the viewing direction.

Images