← Back to branch summary

~mmach/netext73/mesa-ryzen

« back to all changes in this revision

Viewing changes to src/asahi/compiler/agx_opcodes.py

  • Committer: mmach
  • Date: 2023-11-02 21:31:35 UTC
  • Revision ID: netbit73@gmail.com-20231102213135-18d4tzh7tj0uz752
2023-11-02 22:11:57

Show diffs side-by-side

added added

removed removed

Lines of Context:
src/asahi/compiler/agx_opcodes.py
108
108
NEST = immediate("nest")
109
109
TARGET = immediate("target", "agx_block *")
110
110
ZS = immediate("zs")
111
 
PERSPECTIVE = immediate("perspective", "bool")
112
111
SR = enum("sr", {
113
112
   0:  'threadgroup_position_in_grid.x',
114
113
   1:  'threadgroup_position_in_grid.y',
119
118
   8:  'dispatch_threads_per_threadgroup.x',
120
119
   9:  'dispatch_threads_per_threadgroup.y',
121
120
   10: 'dispatch_threads_per_threadgroup.z',
 
121
   20: 'core_index',
 
122
   21: 'vm_slot',
122
123
   48: 'thread_position_in_threadgroup.x',
123
124
   49: 'thread_position_in_threadgroup.y',
124
125
   50: 'thread_position_in_threadgroup.z',
132
133
   80: 'thread_position_in_grid.x',
133
134
   81: 'thread_position_in_grid.y',
134
135
   82: 'thread_position_in_grid.z',
 
136
   124: 'input_sample_mask',
 
137
   144: 'opfifo_cmd',
 
138
   146: 'opfifo_data_l',
 
139
   147: 'opfifo_data_h',
135
140
})
136
141
 
137
142
ATOMIC_OPC = enum("atomic_opc", {
148
153
        10: 'xor',
149
154
})
150
155
 
 
156
INTERPOLATION = enum("interpolation", {
 
157
    0: 'center',
 
158
    1: 'sample',
 
159
    2: 'centroid',
 
160
    # We translate sample -> sample_register at pack time for simplicity
 
161
    3: 'sample_register',
 
162
})
 
163
 
151
164
FUNOP = lambda x: (x << 28)
152
165
FUNOP_MASK = FUNOP((1 << 14) - 1)
153
166
 
287
300
 
288
301
op("get_sr", (0x72, 0x7F | L, 4, _), dests = 1, imms = [SR])
289
302
 
290
 
op("sample_mask", (0x7fc1, 0xffff, 6, _), dests = 0, srcs = 1, can_eliminate = False)
 
303
op("sample_mask", (0x7fc1, 0xffff, 6, _), dests = 0, srcs = 2, can_eliminate = False)
291
304
 
292
305
# Sources: sample mask, combined depth/stencil
293
306
op("zs_emit", (0x41, 0xFF | L, 4, _), dests = 0, srcs = 2,
322
335
 
323
336
op("bitop", (0x7E, 0x7F, 6, _), srcs = 2, imms = [TRUTH_TABLE])
324
337
op("convert", (0x3E | L, 0x7F | L | (0x3 << 38), 6, _), srcs = 2, imms = [ROUND]) 
325
 
op("iter", (0x21, 0xBF, 8, _), srcs = 2, imms = [CHANNELS, PERSPECTIVE])
 
338
 
 
339
# Sources are the coeffient register and the sample index (if applicable)
 
340
op("iter", (0x21, 0xBF, 8, _), srcs = 2, imms = [CHANNELS, INTERPOLATION])
 
341
 
 
342
# Sources are the coeffient register for the varying, the coefficient register
 
343
# for W, and the sample index (if applicable)
 
344
op("iterproj", (0x21, 0xBF, 8, _), srcs = 3, imms = [CHANNELS, INTERPOLATION])
 
345
 
326
346
op("ldcf", (0xA1, 0xBF, 8, _), srcs = 1, imms = [CHANNELS])
327
347
op("st_vary", None, dests = 0, srcs = 2, can_eliminate = False)
328
348
op("no_varyings", (0x80000051, 0xFFFFFFFF, 4, _), dests = 0, can_eliminate = False)