Skip to content

Instantly share code, notes, and snippets.

@qpwo
Last active July 28, 2026 11:03
Show Gist options
  • Select an option

  • Save qpwo/e13358c1d7dba2f664027e5e02b9bd63 to your computer and use it in GitHub Desktop.

Select an option

Save qpwo/e13358c1d7dba2f664027e5e02b9bd63 to your computer and use it in GitHub Desktop.
safe assembly thingy
; WOWNEAT/3. Intel AMD64 plus bnd. No language runtime operations.
; Full-width pointer moves preserve hidden authority. Integer writes clear it.
; Every imported entry is verifier-described in section .wowneat.abi.
; Linux x86-64 SysV register order is used for direct calls.
; Verifier must track allocation identity per bnd-tagged pointer through all
; full-width moves and memory loads. When rt.free is called on a pointer, the
; verifier marks that allocation as freed. Any subsequent load or store through
; any pointer whose provenance traces to a freed allocation is a compile-time
; error. This closes use-after-free: stale bnd bounds cannot match a live
; allocation because the verifier rejects the dereference before code emission.
; Integer-cleared authority (byte copies, arithmetic) cannot reconstruct bnd
; tags, so there is no escape path from freed-status tracking.
bits 64
default rel
extern rt.alloc
extern rt.free
extern os.exit
extern os.open_ro
extern os.open_dir
extern os.open_rw_create
extern os.open_rw_resume
extern os.close
extern os.read
extern os.pread
extern os.write
extern os.pwrite
extern os.seek_end
extern os.truncate
extern os.fsync
extern os.rename
extern os.unlink
extern os.mkdir
extern os.map_ro
extern os.unmap
extern os.clock_ns
extern os.sleep_ns
extern os.cpu_count
extern os.thread_start
extern os.thread_join
extern os.wait32
extern os.wake32
extern os.socket_tcp
extern os.bind_listen
extern os.accept
extern os.recv
extern os.send
extern os.shutdown
extern tls.open
extern tls.read
extern tls.write
extern tls.close
extern gpu.count
extern gpu.open
extern gpu.close
extern gpu.alloc
extern gpu.free
extern gpu.copy_h2d
extern gpu.copy_d2h
extern gpu.copy_peer
extern gpu.peer_ok
extern gpu.kernel
extern gpu.launch
extern gpu.sync
extern gpu.event_new
extern gpu.event_record
extern gpu.event_wait
extern gpu.event_free
O_RDONLY equ 0
O_WRONLY equ 1
O_RDWR equ 2
EINTR equ -4
EAGAIN equ -11
EEXIST equ -17
HTTP_HEAD_MAX equ 32768
HTTP_REQ_MAX equ 65536
HTTP_URL_MAX equ 2048
TLS_PORT equ 443
PAGE equ 4096
section .rodata
s_usage db "mode: all|download|train|eval|serve root [port]",10
s_usage_n equ $-s_usage
s_all db "all",0
s_download db "download",0
s_train db "train",0
s_eval db "eval",0
s_serve db "serve",0
s_http11 db "HTTP/1.1 ",0
s_crlf db 13,10,0
s_crlf2 db 13,10,13,10,0
s_get db "GET ",0
s_post db "POST ",0
s_content_length db "content-length",0
s_transfer_encoding db "transfer-encoding",0
s_connection db "connection",0
s_location db "location",0
s_chunked db "chunked",0
s_close db "close",0
s_identity db "identity",0
s_host db "host",0
s_range db "range",0
s_accept db "accept",0
s_ct db "content-type",0
s_ok db "HTTP/1.1 200 OK",13,10,"Connection: close",13,10,0
s_bad db "HTTP/1.1 400 Bad Request",13,10,"Connection: close",13,10,"Content-Length: 0",13,10,13,10
s_bad_n equ $-s_bad
s_notfound db "HTTP/1.1 404 Not Found",13,10,"Connection: close",13,10,"Content-Length: 0",13,10,13,10
s_notfound_n equ $-s_notfound
s_busy db "HTTP/1.1 503 Busy",13,10,"Connection: close",13,10,"Content-Length: 0",13,10,13,10
s_busy_n equ $-s_busy
hex_lower db "0123456789abcdef"
size hex_lower,16
section .text
fatal:
; rdi=message authority, rsi=bytes, edx=exit code
push rdx
mov rdx,rsi
mov rsi,rdi
mov edi,2
call os.write
pop rdi
call os.exit
ud2
usage:
lea rdi,[rel s_usage]
mov rsi,s_usage_n
mov edx,64
jmp fatal
alloc:
; rdi=bytes; rax=exact zeroed owning authority
push rdi
call rt.alloc
pop rcx
test rax,rax
jz oom
bnd rax,rcx
ret
oom:
lea rdi,[rel s_oom]
mov rsi,s_oom_n
mov edx,70
jmp fatal
free:
test rdi,rdi
jz free_ret
call rt.free
free_ret:
ret
streq:
; rdi,rsi zero-terminated; eax=1 equal
streq_loop:
mov al,byte [rdi]
mov dl,byte [rsi]
cmp al,dl
jne streq_no
test al,al
jz streq_yes
lea rdi,[rdi+1]
lea rsi,[rsi+1]
jmp streq_loop
streq_yes:
mov eax,1
ret
streq_no:
xor eax,eax
ret
strlen:
; rdi; rax bytes excluding zero
mov rax,rdi
strlen_loop:
cmp byte [rax],0
je strlen_done
lea rax,[rax+1]
jmp strlen_loop
strlen_done:
sub rax,rdi
ret
memzero:
; rdi authority, rsi bytes
xor eax,eax
mov rcx,rsi
rep stosb
ret
memcpy_data:
; rdi dst, rsi src, rdx bytes; byte copy intentionally clears pointer tags
mov rcx,rdx
rep movsb
ret
memcmp:
; rdi,rsi,rdx; eax=-1/0/1
xor ecx,ecx
memcmp_loop:
cmp rcx,rdx
je memcmp_eq
mov al,byte [rdi+rcx]
mov r8b,byte [rsi+rcx]
cmp al,r8b
jb memcmp_lt
ja memcmp_gt
inc rcx
jmp memcmp_loop
memcmp_eq:
xor eax,eax
ret
memcmp_lt:
mov eax,-1
ret
memcmp_gt:
mov eax,1
ret
strieq_n:
; rdi,rsi,rdx exact length; eax=1
xor ecx,ecx
strieq_n_loop:
cmp rcx,rdx
je strieq_n_yes
mov al,byte [rdi+rcx]
mov r8b,byte [rsi+rcx]
cmp al,'A'
jb strieq_n_skip1
cmp al,'Z'
ja strieq_n_skip1
add al,32
strieq_n_skip1:
cmp r8b,'A'
jb strieq_n_skip2
cmp r8b,'Z'
ja strieq_n_skip2
add r8b,32
strieq_n_skip2:
cmp al,r8b
jne strieq_n_no
inc rcx
jmp strieq_n_loop
strieq_n_yes:
mov eax,1
ret
strieq_n_no:
xor eax,eax
ret
parse_u64:
; rdi bytes, rsi count; rax value, edx=0 success
xor eax,eax
xor ecx,ecx
test rsi,rsi
jz parse_u64_bad
parse_u64_loop:
cmp rcx,rsi
je parse_u64_ok
movzx r8d,byte [rdi+rcx]
sub r8d,'0'
cmp r8d,9
ja parse_u64_bad
mov r9,1844674407370955161
cmp rax,r9
ja parse_u64_bad
imul rax,rax,10
add rax,r8
jc parse_u64_bad
inc rcx
jmp parse_u64_loop
parse_u64_ok:
xor edx,edx
ret
parse_u64_bad:
mov edx,1
ret
u64_dec:
; rdi dst >=20, rsi value; rax count
push rbx
push r12
lea rsp,[rsp-32]
mov rbx,rdi
lea r12,[rsp+32]
mov rax,rsi
xor ecx,ecx
u64_dec_loop:
xor edx,edx
mov r9,10
div r9
add dl,'0'
dec r12
mov byte [r12],dl
inc ecx
test rax,rax
jnz u64_dec_loop
mov rdi,rbx
mov rsi,r12
mov edx,ecx
push rcx
call memcpy_data
pop rax
lea rsp,[rsp+32]
pop r12
pop rbx
ret
file_write_all:
; edi fd, rsi buffer, rdx bytes; eax=0 or negative
push rbx
mov ebx,edi
file_write_all_loop:
test rdx,rdx
jz file_write_all_ok
mov edi,ebx
call os.write
cmp rax,EINTR
je file_write_all_loop
test rax,rax
jle file_write_all_ret
add rsi,rax
sub rdx,rax
jmp file_write_all_loop
file_write_all_ok:
xor eax,eax
file_write_all_ret:
pop rbx
ret
file_read_exact:
; edi fd,rsi dst,rdx count; eax=0 success
push rbx
mov ebx,edi
file_read_exact_loop:
test rdx,rdx
jz file_read_exact_ok
mov edi,ebx
call os.read
cmp rax,EINTR
je file_read_exact_loop
test rax,rax
jle file_read_exact_bad
add rsi,rax
sub rdx,rax
jmp file_read_exact_loop
file_read_exact_ok:
xor eax,eax
pop rbx
ret
file_read_exact_bad:
mov eax,-1
pop rbx
ret
file_pread_exact:
; edi fd,rsi dst,rdx count,rcx offset
push rbx
push r12
mov ebx,edi
mov r12,rcx
file_pread_exact_loop:
test rdx,rdx
jz file_pread_exact_ok
mov edi,ebx
mov rcx,r12
call os.pread
cmp rax,EINTR
je file_pread_exact_loop
test rax,rax
jle file_pread_exact_bad
add rsi,rax
add r12,rax
sub rdx,rax
jmp file_pread_exact_loop
file_pread_exact_ok:
xor eax,eax
pop r12
pop rbx
ret
file_pread_exact_bad:
mov eax,-1
pop r12
pop rbx
ret
file_map_exact:
; rdi path, rsi expected bytes; rax mapping, rdx bytes
push rbx
mov rbx,rsi
call os.open_ro
test eax,eax
js file_map_exact_bad
mov edi,eax
push rdi
call os.seek_end
cmp rax,rbx
jne file_map_exact_close_bad
pop rdi
mov rsi,rbx
call os.map_ro
mov rdx,rbx
pop rbx
ret
file_map_exact_close_bad:
pop rdi
call os.close
file_map_exact_bad:
xor eax,eax
xor edx,edx
pop rbx
ret
s_oom db "out of memory",10
s_oom_n equ $-s_oom
; SHA-256 context: h[8]@0, total_bytes@32, used@40, block[64]@48, 112 bytes.
section .rodata align=64
sha256_k:
dd 0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5
dd 0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174
dd 0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da
dd 0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967
dd 0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85
dd 0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070
dd 0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3
dd 0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2
size sha256_k,256
section .text
sha256_init:
; rdi context >=112
mov dword [rdi+0],0x6a09e667
mov dword [rdi+4],0xbb67ae85
mov dword [rdi+8],0x3c6ef372
mov dword [rdi+12],0xa54ff53a
mov dword [rdi+16],0x510e527f
mov dword [rdi+20],0x9b05688c
mov dword [rdi+24],0x1f83d9ab
mov dword [rdi+28],0x5be0cd19
mov qword [rdi+32],0
mov qword [rdi+40],0
ret
sha256_update:
; rdi context,rsi data,rdx bytes
push rbx
push r12
push r13
mov rbx,rdi
mov r12,rsi
mov r13,rdx
add qword [rbx+32],r13
sha256_update_loop:
test r13,r13
jz sha256_update_done
mov rax,64
sub rax,qword [rbx+40]
cmp rax,r13
cmova rax,r13
lea rdi,[rbx+48]
add rdi,qword [rbx+40]
mov rsi,r12
mov rdx,rax
push rax
call memcpy_data
pop rax
add qword [rbx+40],rax
add r12,rax
sub r13,rax
cmp qword [rbx+40],64
jne sha256_update_loop
mov rdi,rbx
lea rsi,[rbx+48]
call sha256_compress
mov qword [rbx+40],0
jmp sha256_update_loop
sha256_update_done:
pop r13
pop r12
pop rbx
ret
sha256_final:
; rdi context,rsi output32
push rbx
push r12
mov rbx,rdi
mov r12,rsi
mov rax,qword [rbx+32]
shl rax,3
push rax
mov rcx,qword [rbx+40]
lea rdi,[rbx+48+rcx]
mov byte [rdi],0x80
inc rcx
cmp rcx,56
jbe sha256_final_pad56
mov rdx,64
sub rdx,rcx
lea rdi,[rbx+48+rcx]
mov rsi,rdx
call memzero
mov rdi,rbx
lea rsi,[rbx+48]
call sha256_compress
xor ecx,ecx
sha256_final_pad56:
mov rdx,56
sub rdx,rcx
lea rdi,[rbx+48+rcx]
mov rsi,rdx
call memzero
pop rax
bswap rax
mov qword [rbx+48+56],rax
mov rdi,rbx
lea rsi,[rbx+48]
call sha256_compress
xor ecx,ecx
sha256_final_out:
mov eax,dword [rbx+rcx*4]
bswap eax
mov dword [r12+rcx*4],eax
inc ecx
cmp ecx,8
jb sha256_final_out
pop r12
pop rbx
ret
sha256_compress:
; rdi=context,rsi=64-byte block
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-320]
mov rbx,rdi
mov r12,rsi
xor ecx,ecx
sha256_w0:
mov eax,dword [r12+rcx*4]
bswap eax
mov dword [rsp+rcx*4],eax
inc ecx
cmp ecx,16
jb sha256_w0
sha256_wx:
mov eax,dword [rsp+rcx*4-60]
mov edx,eax
ror edx,7
mov r8d,eax
ror r8d,18
xor edx,r8d
shr eax,3
xor edx,eax
mov eax,dword [rsp+rcx*4-8]
mov r8d,eax
ror r8d,17
mov r9d,eax
ror r9d,19
xor r8d,r9d
shr eax,10
xor r8d,eax
add edx,dword [rsp+rcx*4-64]
add edx,dword [rsp+rcx*4-28]
add edx,r8d
mov dword [rsp+rcx*4],edx
inc ecx
cmp ecx,64
jb sha256_wx
xor ecx,ecx
sha256_state_copy:
mov eax,dword [rbx+rcx*4]
mov dword [rsp+256+rcx*4],eax
inc ecx
cmp ecx,8
jb sha256_state_copy
xor ecx,ecx
lea r15,[rel sha256_k]
sha256_round:
; state a..h at 256..284
mov eax,dword [rsp+256+4*4]
mov edx,eax
ror edx,6
mov r8d,eax
ror r8d,11
xor edx,r8d
mov r8d,eax
ror r8d,25
xor edx,r8d
mov r8d,dword [rsp+256+5*4]
mov r9d,dword [rsp+256+6*4]
xor r9d,r8d
and r9d,eax
xor r9d,dword [rsp+256+6*4]
add edx,dword [rsp+256+7*4]
add edx,r9d
add edx,dword [r15+rcx*4]
add edx,dword [rsp+rcx*4]
mov eax,dword [rsp+256+0*4]
mov r8d,eax
ror r8d,2
mov r9d,eax
ror r9d,13
xor r8d,r9d
mov r9d,eax
ror r9d,22
xor r8d,r9d
mov r9d,dword [rsp+256+1*4]
mov r10d,dword [rsp+256+2*4]
mov r11d,eax
and r11d,r9d
mov r13d,eax
and r13d,r10d
xor r11d,r13d
and r9d,r10d
xor r11d,r9d
add r8d,r11d
mov eax,dword [rsp+256+6*4]
mov dword [rsp+256+7*4],eax
mov eax,dword [rsp+256+5*4]
mov dword [rsp+256+6*4],eax
mov eax,dword [rsp+256+4*4]
mov dword [rsp+256+5*4],eax
mov eax,dword [rsp+256+3*4]
add eax,edx
mov dword [rsp+256+4*4],eax
mov eax,dword [rsp+256+2*4]
mov dword [rsp+256+3*4],eax
mov eax,dword [rsp+256+1*4]
mov dword [rsp+256+2*4],eax
mov eax,dword [rsp+256+0*4]
mov dword [rsp+256+1*4],eax
add edx,r8d
mov dword [rsp+256+0*4],edx
inc ecx
cmp ecx,64
jb sha256_round
xor ecx,ecx
sha256_state_add:
mov eax,dword [rsp+256+rcx*4]
add dword [rbx+rcx*4],eax
inc ecx
cmp ecx,8
jb sha256_state_add
lea rsp,[rsp+320]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
sha256_file:
; rdi path,rsi output32; eax=0 success
push rbx
push r12
push r13
mov r12,rsi
call os.open_ro
test eax,eax
js sha256_file_bad
mov ebx,eax
mov rdi,112
call alloc
mov r13,rax
mov rdi,r13
call sha256_init
mov rdi,65536
call alloc
push rax
sha256_file_read:
mov edi,ebx
mov rsi,qword [rsp]
mov rdx,65536
call os.read
cmp rax,EINTR
je sha256_file_read
test rax,rax
js sha256_file_cleanup_bad
jz sha256_file_finish
mov rdi,r13
mov rsi,qword [rsp]
mov rdx,rax
call sha256_update
jmp sha256_file_read
sha256_file_finish:
mov rdi,r13
mov rsi,r12
call sha256_final
mov edi,ebx
call os.close
pop rdi
call free
mov rdi,r13
call free
xor eax,eax
pop r13
pop r12
pop rbx
ret
sha256_file_cleanup_bad:
mov edi,ebx
call os.close
pop rdi
call free
mov rdi,r13
call free
sha256_file_bad:
mov eax,-1
pop r13
pop r12
pop rbx
ret
; HTTP response state: status@0, length@8, flags@16, range_start@24.
; flags bit0 length seen, bit1 chunked, bit2 range seen.
http_append:
; rdi write cursor,rsi source,rdx count; rax new cursor
mov r8,rdi
call memcpy_data
lea rax,[r8+rdx]
ret
http_build_get:
; rdi buffer,rsi capacity,rdx host,rcx path,r8 range offset
; rax bytes, edx=0 success
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-32]
mov rbx,rdi
lea r12,[rdi+rsi]
mov r13,rdx
mov r14,rcx
mov r15,r8
mov rdi,r13
call strlen
mov qword [rsp+0],rax
mov rdi,r14
call strlen
mov qword [rsp+8],rax
mov rax,http_get_prefix_n+http_get_mid_n+http_get_tail_n+http_end_n
add rax,qword [rsp+0]
add rax,qword [rsp+8]
test r15,r15
jz http_build_get_size_ok
add rax,http_range_prefix_n+http_range_tail_n+20
http_build_get_size_ok:
lea rcx,[rbx+rax]
cmp rcx,r12
ja http_build_get_bad
mov rdi,rbx
lea rsi,[rel http_get_prefix]
mov rdx,http_get_prefix_n
call http_append
mov rdi,rax
mov rsi,r14
mov rdx,qword [rsp+8]
call http_append
mov rdi,rax
lea rsi,[rel http_get_mid]
mov rdx,http_get_mid_n
call http_append
mov rdi,rax
mov rsi,r13
mov rdx,qword [rsp+0]
call http_append
mov rdi,rax
lea rsi,[rel http_get_tail]
mov rdx,http_get_tail_n
call http_append
mov rdi,rax
test r15,r15
jz http_build_get_end
lea rsi,[rel http_range_prefix]
mov rdx,http_range_prefix_n
call http_append
mov rdi,rax
mov rsi,r15
call u64_dec
add rdi,rax
lea rsi,[rel http_range_tail]
mov rdx,http_range_tail_n
call http_append
mov rdi,rax
http_build_get_end:
lea rsi,[rel http_end]
mov rdx,http_end_n
call http_append
sub rax,rbx
xor edx,edx
jmp http_build_get_ret
http_build_get_bad:
xor eax,eax
mov edx,1
http_build_get_ret:
lea rsp,[rsp+32]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
section .rodata
http_get_prefix db "GET "
http_get_prefix_n equ $-http_get_prefix
http_get_mid db " HTTP/1.1",13,10,"Host: "
http_get_mid_n equ $-http_get_mid
http_get_tail db 13,10,"User-Agent: wowneat/3",13,10,"Accept: application/octet-stream",13,10,"Accept-Encoding: identity",13,10,"Connection: close",13,10
http_get_tail_n equ $-http_get_tail
http_range_prefix db "Range: bytes="
http_range_prefix_n equ $-http_range_prefix
http_range_tail db "-",13,10
http_range_tail_n equ $-http_range_tail
http_end db 13,10
http_end_n equ $-http_end
section .text
find_crlf:
; rdi start,rsi bytes; rax offset or -1
xor eax,eax
find_crlf_loop:
lea rcx,[rax+1]
cmp rcx,rsi
jae find_crlf_no
cmp byte [rdi+rax],13
jne find_crlf_next
cmp byte [rdi+rax+1],10
je find_crlf_yes
find_crlf_next:
inc rax
jmp find_crlf_loop
find_crlf_no:
mov rax,-1
find_crlf_yes:
ret
tls_read_headers:
; rdi stream,rsi buffer,rdx capacity; rax total,rdx header_end; negative on error
push rbx
push r12
push r13
push r14
mov rbx,rdi
mov r12,rsi
mov r13,rdx
xor r14d,r14d
tls_read_headers_read:
cmp r14,r13
jae tls_read_headers_bad
mov rdi,rbx
lea rsi,[r12+r14]
mov rdx,r13
sub rdx,r14
call tls.read
cmp rax,EINTR
je tls_read_headers_read
test rax,rax
jle tls_read_headers_bad
mov rcx,r14
add r14,rax
cmp rcx,3
jae tls_read_headers_scan
xor ecx,ecx
tls_read_headers_scan:
lea rax,[rcx+3]
cmp rax,r14
jae tls_read_headers_read
cmp dword [r12+rcx],0x0a0d0a0d
je tls_read_headers_found
inc rcx
jmp tls_read_headers_scan
tls_read_headers_found:
lea rdx,[rcx+4]
mov rax,r14
pop r14
pop r13
pop r12
pop rbx
ret
tls_read_headers_bad:
mov rax,-1
xor edx,edx
pop r14
pop r13
pop r12
pop rbx
ret
http_trim_value:
; rdi value start,rsi bytes; rax start,rdx bytes
mov rax,rdi
mov rdx,rsi
http_trim_left:
test rdx,rdx
jz http_trim_done
mov cl,byte [rax]
cmp cl,' '
je http_trim_left_one
cmp cl,9
jne http_trim_right
http_trim_left_one:
inc rax
dec rdx
jmp http_trim_left
http_trim_right:
test rdx,rdx
jz http_trim_done
mov cl,byte [rax+rdx-1]
cmp cl,' '
je http_trim_right_one
cmp cl,9
jne http_trim_done
http_trim_right_one:
dec rdx
jmp http_trim_right
http_trim_done:
ret
http_parse_content_range:
; rdi value,rsi len; rax start,edx=0 success
cmp rsi,8
jb http_parse_content_range_bad
cmp dword [rdi],0x65747962
jne http_parse_content_range_bad
cmp byte [rdi+4],'s'
jne http_parse_content_range_bad
cmp byte [rdi+5],' '
jne http_parse_content_range_bad
lea rdi,[rdi+6]
sub rsi,6
xor ecx,ecx
http_parse_content_range_dash:
cmp rcx,rsi
jae http_parse_content_range_bad
cmp byte [rdi+rcx],'-'
je http_parse_content_range_number
inc rcx
jmp http_parse_content_range_dash
http_parse_content_range_number:
test rcx,rcx
jz http_parse_content_range_bad
mov rsi,rcx
call parse_u64
ret
http_parse_content_range_bad:
xor eax,eax
mov edx,1
ret
http_parse_response:
; rdi buffer,rsi header_end,rdx state32; eax=0 success
push rbx
push r12
push r13
push r14
push r15
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov qword [r13+0],0
mov qword [r13+8],0
mov qword [r13+16],0
mov qword [r13+24],0
mov rdi,rbx
mov rsi,r12
call find_crlf
cmp rax,12
jb http_parse_response_bad
mov r14,rax
lea rdi,[rbx]
lea rsi,[rel s_http11]
mov rdx,9
call memcmp
test eax,eax
jnz http_parse_response_bad
movzx eax,byte [rbx+9]
sub eax,'0'
cmp eax,9
ja http_parse_response_bad
imul eax,eax,100
movzx ecx,byte [rbx+10]
sub ecx,'0'
cmp ecx,9
ja http_parse_response_bad
imul ecx,ecx,10
add eax,ecx
movzx ecx,byte [rbx+11]
sub ecx,'0'
cmp ecx,9
ja http_parse_response_bad
add eax,ecx
mov qword [r13+0],rax
lea r14,[r14+2]
http_parse_response_line:
lea rax,[r14+2]
cmp rax,r12
ja http_parse_response_bad
cmp word [rbx+r14],0x0a0d
je http_parse_response_done
lea rdi,[rbx+r14]
mov rsi,r12
sub rsi,r14
call find_crlf
test rax,rax
jle http_parse_response_bad
mov r15,rax
xor ecx,ecx
http_parse_response_colon:
cmp rcx,r15
jae http_parse_response_bad
cmp byte [rbx+r14+rcx],':'
je http_parse_response_have_colon
inc rcx
jmp http_parse_response_colon
http_parse_response_have_colon:
mov r8,rcx
lea rdi,[rbx+r14+rcx+1]
mov rsi,r15
sub rsi,rcx
sub rsi,1
call http_trim_value
mov r9,rax
mov r10,rdx
lea rdi,[rbx+r14]
lea rsi,[rel s_content_length]
mov rdx,14
cmp r8,14
jne http_parse_response_te
call strieq_n
test eax,eax
jz http_parse_response_te
test qword [r13+16],1
jnz http_parse_response_bad
mov rdi,r9
mov rsi,r10
call parse_u64
test edx,edx
jnz http_parse_response_bad
mov qword [r13+8],rax
or qword [r13+16],1
jmp http_parse_response_next
http_parse_response_te:
lea rdi,[rbx+r14]
lea rsi,[rel s_transfer_encoding]
mov rdx,17
cmp r8,17
jne http_parse_response_range
call strieq_n
test eax,eax
jz http_parse_response_range
mov rdi,r9
lea rsi,[rel s_identity]
mov rdx,8
cmp r10,8
jne http_parse_response_te_chunked
call strieq_n
test eax,eax
jnz http_parse_response_next
http_parse_response_te_chunked:
mov rdi,r9
lea rsi,[rel s_chunked]
mov rdx,7
cmp r10,7
jne http_parse_response_bad
call strieq_n
test eax,eax
jz http_parse_response_bad
or qword [r13+16],2
jmp http_parse_response_next
http_parse_response_range:
lea rdi,[rbx+r14]
lea rsi,[rel s_content_range]
mov rdx,13
cmp r8,13
jne http_parse_response_next
call strieq_n
test eax,eax
jz http_parse_response_next
mov rdi,r9
mov rsi,r10
call http_parse_content_range
test edx,edx
jnz http_parse_response_bad
mov qword [r13+24],rax
or qword [r13+16],4
http_parse_response_next:
add r14,r15
add r14,2
jmp http_parse_response_line
http_parse_response_done:
mov rax,qword [r13+16]
test rax,2
jnz http_parse_response_bad
test rax,1
jz http_parse_response_bad
xor eax,eax
jmp http_parse_response_ret
http_parse_response_bad:
mov eax,-1
http_parse_response_ret:
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
section .rodata
s_content_range db "content-range",0
section .text
tls_write_all:
; rdi stream,rsi data,rdx bytes; eax=0 success
push rbx
mov rbx,rdi
tls_write_all_loop:
test rdx,rdx
jz tls_write_all_ok
mov rdi,rbx
call tls.write
cmp rax,EINTR
je tls_write_all_loop
test rax,rax
jle tls_write_all_bad
add rsi,rax
sub rdx,rax
jmp tls_write_all_loop
tls_write_all_ok:
xor eax,eax
pop rbx
ret
tls_write_all_bad:
mov eax,-1
pop rbx
ret
https_download_pinned:
; We tailcall https_download_once but force rcx (part) to equal rdx (final)
; This writes directly to the final path with resume, skipping a .part file.
mov rcx,rdx
jmp https_download_once
https_download_once:
; same args as https_download_pinned; writes part, verifies byte count only
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-160]
mov qword [rsp+0],rdi
mov qword [rsp+8],rsi
mov qword [rsp+16],rdx
mov qword [rsp+24],rcx
mov qword [rsp+32],r8
mov qword [rsp+40],r9
mov rdi,rcx
call os.open_rw_resume
test eax,eax
js https_download_once_bad
mov ebx,eax
mov edi,ebx
call os.seek_end
test rax,rax
js https_download_once_close_bad
mov r12,rax
cmp r12,qword [rsp+32]
jbe https_download_once_size_ok
mov edi,ebx
xor esi,esi
call os.truncate
test eax,eax
js https_download_once_close_bad
xor r12d,r12d
https_download_once_size_ok:
mov rdi,qword [rsp+0]
mov esi,TLS_PORT
call tls.open
test rax,rax
jz https_download_once_close_bad
mov r13,rax
mov rdi,65536
call alloc
mov r14,rax
mov rdi,r14
mov rsi,65536
mov rdx,qword [rsp+0]
mov rcx,qword [rsp+8]
mov r8,r12
call http_build_get
test edx,edx
jnz https_download_once_tls_bad
mov r15,rax
mov rdi,r13
mov rsi,r14
mov rdx,r15
call tls_write_all
test eax,eax
jnz https_download_once_tls_bad
mov rdi,r13
mov rsi,r14
mov rdx,65536
call tls_read_headers
test rax,rax
js https_download_once_tls_bad
mov qword [rsp+48],rax
mov qword [rsp+56],rdx
lea rdx,[rsp+80]
mov rdi,r14
mov rsi,qword [rsp+56]
call http_parse_response
test eax,eax
jnz https_download_once_tls_bad
mov rax,qword [rsp+80]
test r12,r12
jz https_download_once_status_zero
cmp rax,206
jne https_download_once_restart_zero
test qword [rsp+96],4
jz https_download_once_tls_bad
cmp qword [rsp+104],r12
jne https_download_once_tls_bad
jmp https_download_once_length
https_download_once_status_zero:
cmp rax,200
je https_download_once_length
cmp rax,206
jne https_download_once_tls_bad
test qword [rsp+96],4
jz https_download_once_tls_bad
cmp qword [rsp+104],0
jne https_download_once_tls_bad
https_download_once_length:
mov rax,qword [rsp+32]
sub rax,r12
cmp qword [rsp+88],rax
jne https_download_once_tls_bad
mov r15,qword [rsp+48]
sub r15,qword [rsp+56]
test r15,r15
jz https_download_once_body_loop
cmp r15,qword [rsp+88]
ja https_download_once_tls_bad
mov edi,ebx
mov rax,qword [rsp+56]
lea rsi,[r14+rax]
mov rdx,r15
call file_write_all
test eax,eax
jnz https_download_once_tls_bad
add r12,r15
https_download_once_body_loop:
cmp r12,qword [rsp+32]
je https_download_once_body_done
mov rdi,r13
mov rsi,r14
mov rdx,65536
call tls.read
cmp rax,EINTR
je https_download_once_body_loop
test rax,rax
jle https_download_once_tls_bad
mov rcx,qword [rsp+32]
sub rcx,r12
cmp rax,rcx
ja https_download_once_tls_bad
mov edi,ebx
mov rsi,r14
mov rdx,rax
push rax
call file_write_all
pop rcx
test eax,eax
jnz https_download_once_tls_bad
add r12,rcx
jmp https_download_once_body_loop
https_download_once_body_done:
mov rdi,r13
call tls.close
mov rdi,r14
call free
mov edi,ebx
call os.fsync
mov edi,ebx
call os.close
xor eax,eax
jmp https_download_once_ret
https_download_once_restart_zero:
mov rdi,r13
call tls.close
mov rdi,r14
call free
mov edi,ebx
xor esi,esi
call os.truncate
mov edi,ebx
call os.close
mov eax,-2
jmp https_download_once_ret
https_download_once_tls_bad:
mov rdi,r13
call tls.close
mov rdi,r14
call free
https_download_once_close_bad:
mov edi,ebx
call os.close
https_download_once_bad:
mov eax,-1
https_download_once_ret:
lea rsp,[rsp+160]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
; Request state: method@0, path_off@8, path_len@16, body_off@24, body_len@32.
http_read_request:
; edi fd,rsi buffer,rdx capacity,rcx state; eax=0 success
push rbx
push r12
push r13
push r14
push r15
mov ebx,edi
mov r12,rsi
mov r13,rdx
mov r14,rcx
mov qword [r14+0],0
mov qword [r14+8],0
mov qword [r14+16],0
mov qword [r14+24],0
mov qword [r14+32],0
xor r15d,r15d
http_read_request_more:
cmp r15,r13
jae http_read_request_bad
mov edi,ebx
lea rsi,[r12+r15]
mov rdx,r13
sub rdx,r15
call os.recv
cmp rax,EINTR
je http_read_request_more
test rax,rax
jle http_read_request_bad
mov rcx,r15
add r15,rax
cmp rcx,3
jae http_read_request_scan
xor ecx,ecx
http_read_request_scan:
lea rax,[rcx+3]
cmp rax,r15
jae http_read_request_more
cmp dword [r12+rcx],0x0a0d0a0d
je http_read_request_header
inc rcx
jmp http_read_request_scan
http_read_request_header:
lea r8,[rcx+4]
mov qword [r14+24],r8
mov rdi,r12
mov rsi,r8
mov rdx,r14
call http_parse_request_header
test eax,eax
jnz http_read_request_bad
mov rax,qword [r14+24]
add rax,qword [r14+32]
cmp rax,r13
ja http_read_request_bad
http_read_request_body:
cmp r15,rax
jae http_read_request_ok
mov edi,ebx
lea rsi,[r12+r15]
mov rdx,rax
sub rdx,r15
call os.recv
cmp rax,EINTR
je http_read_request_body
test rax,rax
jle http_read_request_bad
add r15,rax
mov rax,qword [r14+24]
add rax,qword [r14+32]
jmp http_read_request_body
http_read_request_ok:
cmp r15,rax
jne http_read_request_bad
xor eax,eax
jmp http_read_request_ret
http_read_request_bad:
mov eax,-1
http_read_request_ret:
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
http_parse_request_header:
; rdi buffer,rsi header_end,rdx state; eax=0 success
push rbx
push r12
push r13
push r14
push r15
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov rdi,rbx
mov rsi,r12
call find_crlf
test rax,rax
jle http_parse_request_bad
mov r14,rax
cmp r14,14
jb http_parse_request_bad
lea rdi,[rbx]
lea rsi,[rel s_get]
mov rdx,4
call memcmp
test eax,eax
jnz http_parse_request_post
mov qword [r13+0],1
mov r15,4
jmp http_parse_request_path
http_parse_request_post:
lea rdi,[rbx]
lea rsi,[rel s_post]
mov rdx,5
call memcmp
test eax,eax
jnz http_parse_request_bad
mov qword [r13+0],2
mov r15,5
http_parse_request_path:
mov qword [r13+8],r15
mov rcx,r15
http_parse_request_space:
cmp rcx,r14
jae http_parse_request_bad
cmp byte [rbx+rcx],' '
je http_parse_request_version
cmp byte [rbx+rcx],0x20
jb http_parse_request_bad
inc rcx
jmp http_parse_request_space
http_parse_request_version:
mov rax,rcx
sub rax,r15
test rax,rax
jz http_parse_request_bad
mov qword [r13+16],rax
inc rcx
mov rax,r14
sub rax,rcx
cmp rax,8
jne http_parse_request_bad
cmp qword [rbx+rcx],0x312e312f50545448
jne http_parse_request_bad
lea r14,[r14+2]
xor r15d,r15d
http_parse_request_line:
cmp word [rbx+r14],0x0a0d
je http_parse_request_done
lea rdi,[rbx+r14]
mov rsi,r12
sub rsi,r14
call find_crlf
test rax,rax
jle http_parse_request_bad
mov r8,rax
xor ecx,ecx
http_parse_request_colon:
cmp rcx,r8
jae http_parse_request_bad
cmp byte [rbx+r14+rcx],':'
je http_parse_request_have_colon
inc rcx
jmp http_parse_request_colon
http_parse_request_have_colon:
mov r9,rcx
lea rdi,[rbx+r14+rcx+1]
mov rsi,r8
sub rsi,rcx
sub rsi,1
call http_trim_value
mov r10,rax
mov r11,rdx
cmp r9,14
jne http_parse_request_te
lea rdi,[rbx+r14]
lea rsi,[rel s_content_length]
mov rdx,14
call strieq_n
test eax,eax
jz http_parse_request_te
test r15,r15
jnz http_parse_request_bad
mov rdi,r10
mov rsi,r11
call parse_u64
test edx,edx
jnz http_parse_request_bad
mov qword [r13+32],rax
mov r15,1
jmp http_parse_request_next
http_parse_request_te:
cmp r9,17
jne http_parse_request_next
lea rdi,[rbx+r14]
lea rsi,[rel s_transfer_encoding]
mov rdx,17
call strieq_n
test eax,eax
jz http_parse_request_next
jmp http_parse_request_bad
http_parse_request_next:
add r14,r8
add r14,2
jmp http_parse_request_line
http_parse_request_done:
cmp qword [r13+0],2
jne http_parse_request_get_len
; POST must have exactly one Content-Length.
test r15,r15
jz http_parse_request_bad
jmp http_parse_request_success
http_parse_request_get_len:
cmp qword [r13+32],0
jne http_parse_request_bad
http_parse_request_success:
xor eax,eax
jmp http_parse_request_ret
http_parse_request_bad:
mov eax,-1
http_parse_request_ret:
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
socket_send_all:
; edi fd,rsi data,rdx bytes
push rbx
mov ebx,edi
socket_send_all_loop:
test rdx,rdx
jz socket_send_all_ok
mov edi,ebx
call os.send
cmp rax,EINTR
je socket_send_all_loop
test rax,rax
jle socket_send_all_bad
add rsi,rax
sub rdx,rax
jmp socket_send_all_loop
socket_send_all_ok:
xor eax,eax
pop rbx
ret
socket_send_all_bad:
mov eax,-1
pop rbx
ret
http_send_fixed:
; edi fd,rsi content-type z,rdx body,rcx body bytes
push rbx
push r12
push r13
push r14
lea rsp,[rsp-256]
mov ebx,edi
mov r12,rsi
mov r13,rdx
mov r14,rcx
lea rdi,[rsp]
lea rsi,[rel s_ok]
call strlen
mov rdx,rax
call http_append
mov r8,rax
mov rdi,r8
lea rsi,[rel http_ct_prefix]
mov rdx,http_ct_prefix_n
call http_append
mov rdi,rax
mov rsi,r12
push rdi
mov rdi,r12
call strlen
mov rdx,rax
pop rdi
call http_append
mov rdi,rax
lea rsi,[rel http_cl_prefix]
mov rdx,http_cl_prefix_n
call http_append
mov rdi,rax
mov rsi,r14
call u64_dec
add rdi,rax
lea rsi,[rel s_crlf2]
mov rdx,4
call http_append
sub rax,rsp
mov edi,ebx
mov rsi,rsp
mov rdx,rax
call socket_send_all
test eax,eax
jnz http_send_fixed_ret
mov edi,ebx
mov rsi,r13
mov rdx,r14
call socket_send_all
http_send_fixed_ret:
lea rsp,[rsp+256]
pop r14
pop r13
pop r12
pop rbx
ret
section .rodata
http_ct_prefix db "Content-Type: "
http_ct_prefix_n equ $-http_ct_prefix
http_cl_prefix db 13,10,"Content-Length: "
http_cl_prefix_n equ $-http_cl_prefix
ct_json db "application/json",0
section .text
T_CAP equ 0
T_BYTES equ 8
T_D0 equ 16
T_D1 equ 24
T_D2 equ 32
T_D3 equ 40
T_TYPE equ 48
T_FLAGS equ 56
T_SIZE equ 64
DT_U8 equ 1
DT_I16 equ 2
DT_I32 equ 3
DT_F32 equ 4
DT_BF16 equ 5
DT_I8 equ 6
TF_HOST equ 1
TF_GPU equ 2
TF_PINNED equ 4
GA_VALUE equ 0
GA_OFFSET equ 8
GA_BYTES equ 16
GA_KIND equ 24
GA_SIZE equ 32
GA_SPAN_R equ 1
GA_SPAN_W equ 2
GA_SPAN_RW equ 3
GA_U64 equ 4
GA_F32 equ 5
LD_GRID equ 0
LD_BLOCK equ 8
LD_SMEM equ 16
LD_ARGS equ 24
LD_ARGC equ 32
LD_SIZE equ 40
tensor_host_new:
; rdi descriptor,rsi bytes,rdx d0,rcx d1,r8 d2,r9 d3; type at [rsp+8]
push rbx
mov rbx,rdi
mov qword [rbx+T_BYTES],rsi
mov qword [rbx+T_D0],rdx
mov qword [rbx+T_D1],rcx
mov qword [rbx+T_D2],r8
mov qword [rbx+T_D3],r9
mov rdi,rsi
call alloc
mov qword [rbx+T_CAP],rax
mov rax,qword [rsp+16]
mov qword [rbx+T_TYPE],rax
mov qword [rbx+T_FLAGS],TF_HOST
pop rbx
ret
tensor_host_free:
; rdi descriptor
push rbx
mov rbx,rdi
mov rdi,qword [rbx+T_CAP]
call free
mov qword [rbx+T_CAP],0
mov qword [rbx+T_BYTES],0
pop rbx
ret
tensor_gpu_new:
; rdi device,rsi descriptor,rdx bytes,r8 type
push rbx
push r12
mov r12,rdi
mov rbx,rsi
mov qword [rbx+T_BYTES],rdx
mov qword [rbx+T_TYPE],r8
mov qword [rbx+T_FLAGS],TF_GPU
mov rdi,r12
mov rsi,rdx
call gpu.alloc
test rax,rax
jz oom
mov qword [rbx+T_CAP],rax
pop r12
pop rbx
ret
tensor_gpu_free:
; rdi device,rsi descriptor
push rbx
mov rbx,rsi
mov rsi,qword [rbx+T_CAP]
call gpu.free
mov qword [rbx+T_CAP],0
mov qword [rbx+T_BYTES],0
pop rbx
ret
gpu_arg_span:
; rdi record,rsi capability,rdx offset,rcx bytes,r8 kind
mov qword [rdi+GA_VALUE],rsi
mov qword [rdi+GA_OFFSET],rdx
mov qword [rdi+GA_BYTES],rcx
mov qword [rdi+GA_KIND],r8
ret
gpu_arg_u64:
; rdi record,rsi value
mov qword [rdi+GA_VALUE],rsi
mov qword [rdi+GA_OFFSET],0
mov qword [rdi+GA_BYTES],8
mov qword [rdi+GA_KIND],GA_U64
ret
gpu_arg_f32:
; rdi record,esi bits
mov dword [rdi+GA_VALUE],esi
mov dword [rdi+GA_VALUE+4],0
mov qword [rdi+GA_OFFSET],0
mov qword [rdi+GA_BYTES],4
mov qword [rdi+GA_KIND],GA_F32
ret
gpu_pack_xyz:
; edi x,esi y,edx z; rax=x|y<<21|z<<42, each <=2^21-1
mov eax,edi
mov rcx,rsi
shl rcx,21
or rax,rcx
mov rcx,rdx
shl rcx,42
or rax,rcx
ret
gpu_launch_desc:
; rdi device,rsi sealed kernel,rdx launch descriptor
call gpu.launch
ret
f32_exp:
; xmm0 -> xmm0, x87 exact machine operation sequence
lea rsp,[rsp-16]
movss dword [rsp],xmm0
fld dword [rsp]
fldl2e
fmulp st1,st0
fld st0
frndint
fsub st1,st0
fxch st1
f2xm1
fld1
faddp st1,st0
fscale
fstp st1
fstp dword [rsp]
movss xmm0,dword [rsp]
lea rsp,[rsp+16]
ret
f32_cos:
lea rsp,[rsp-16]
movss dword [rsp],xmm0
fld dword [rsp]
fcos
fstp dword [rsp]
movss xmm0,dword [rsp]
lea rsp,[rsp+16]
ret
section .rodata align=16
f_one dd 1.0
f_two dd 2.0
f_half dd 0.5
f_pi dd 3.141592653589793
f_inv_u32 dd 2.3283064365386963e-10
section .text
; cfar.wowneat means CIFAR-10. Dataset records are one label byte plus 3072 planar RGB bytes.
; Model: pre-activation residual CNN, GroupNorm, SiLU, FP32 master/AdamW, BF16 SM89 execution.
; 8 GPUs x 125 valid images in 128 storage slots = global batch 1000. 200 epochs.
CFAR_GPU_MAX equ 8
CFAR_CPU_LOADERS equ 96
CFAR_LOCAL_BATCH equ 128
CFAR_VALID_PER_GPU equ 125
CFAR_GLOBAL_BATCH equ 1000
CFAR_TRAIN_IMAGES equ 50000
CFAR_TEST_IMAGES equ 10000
CFAR_RECORD_BYTES equ 3073
CFAR_BATCH_BYTES equ 30730000
CFAR_IMAGE_BYTES equ 3072
CFAR_CLASSES equ 10
CFAR_EPOCHS equ 200
CFAR_STEPS_PER_EPOCH equ 50
CFAR_TOTAL_STEPS equ 10000
CFAR_CHECKPOINT_PERIOD equ 250
CFAR_WARMUP_STEPS equ 250
CFAR_LABEL_SMOOTH_BITS equ 0x3dcccccd
CFAR_LOSS_SCALE_INIT equ 65536
CFAR_OP_SIZE equ 64
CFAR_CKPT_HEADER equ 4096
CFAR_SERVE_WORKERS equ 32
OP_CONV equ 1
OP_GN equ 2
OP_SILU equ 3
OP_ADD equ 4
OP_GAP equ 5
OP_LINEAR equ 6
OP_XENT equ 7
OD_KIND equ 0
OD_IN0 equ 8
OD_IN1 equ 16
OD_OUT equ 24
OD_PARAM equ 32
OD_AUX equ 40
OD_X0 equ 48
OD_X1 equ 56
section .rodata
cfar_host db "www.cs.toronto.edu",0
cfar_url_path db "/~kriz/cifar-10-binary.tar.gz",0
cfar_archive_suffix db "/data/cifar-10-binary.tar.gz",0
cfar_part_suffix db "/data/cifar-10-binary.tar.gz.part",0
cfar_data_suffix db "/data",0
cfar_ckpt_suffix db "/checkpoint/cfar.ckpt",0
cfar_ckpt_tmp_suffix db "/checkpoint/cfar.ckpt.tmp",0
cfar_checkpoint_dir_suffix db "/checkpoint",0
cfar_batch1_suffix db "/data/cifar-10-batches-bin/data_batch_1.bin",0
cfar_batch2_suffix db "/data/cifar-10-batches-bin/data_batch_2.bin",0
cfar_batch3_suffix db "/data/cifar-10-batches-bin/data_batch_3.bin",0
cfar_batch4_suffix db "/data/cifar-10-batches-bin/data_batch_4.bin",0
cfar_batch5_suffix db "/data/cifar-10-batches-bin/data_batch_5.bin",0
cfar_test_suffix db "/data/cifar-10-batches-bin/test_batch.bin",0
cfar_extract_dir_suffix db "/data/cifar-10-batches-bin",0
cfar_archive_sha db 0xc4,0xa3,0x8c,0x50,0xa1,0xbc,0x5f,0x3a,0x1c,0x55,0x37,0xf2,0x15,0x5a,0xb9,0xd6,0x8f,0x9f,0x25,0xeb,0x1e,0xd8,0xd9,0xdd,0xda,0x3d,0xb2,0x9a,0x59,0xbc,0xa1,0xdd
cfar_archive_bytes dq 170052171
cfar_health_path db "/healthz",0
cfar_predict_path db "/v1/cifar10",0
cfar_health_json db "{\"ok\":true,\"model\":\"cifar10-resgn\"}",10
cfar_health_json_n equ $-cfar_health_json
cfar_err_download db "cifar download failed",10
cfar_err_download_n equ $-cfar_err_download
cfar_err_extract db "cifar archive invalid",10
cfar_err_extract_n equ $-cfar_err_extract
cfar_err_data db "cifar records invalid",10
cfar_err_data_n equ $-cfar_err_data
cfar_err_gpu db "eight usable sm89 devices required",10
cfar_err_gpu_n equ $-cfar_err_gpu
cfar_err_ckpt db "cifar checkpoint invalid",10
cfar_err_ckpt_n equ $-cfar_err_ckpt
cfar_err_train db "cifar training failed",10
cfar_err_train_n equ $-cfar_err_train
cfar_err_eval db "cifar evaluation failed",10
cfar_err_eval_n equ $-cfar_err_eval
cfar_magic db "WOWCFAR3CKPT",0,0,0,0
size cfar_magic,16
cfar_class0 db "airplane",0
cfar_class1 db "automobile",0
cfar_class2 db "bird",0
cfar_class3 db "cat",0
cfar_class4 db "deer",0
cfar_class5 db "dog",0
cfar_class6 db "frog",0
cfar_class7 db "horse",0
cfar_class8 db "ship",0
cfar_class9 db "truck",0
cfar_class_ptrs dq cfar_class0,cfar_class1,cfar_class2,cfar_class3,cfar_class4,cfar_class5,cfar_class6,cfar_class7,cfar_class8,cfar_class9
section .bss align=64
cfar_root resb 4096
cfar_archive_path resb 4096
cfar_part_path resb 4096
cfar_data_path resb 4096
cfar_extract_path resb 4096
cfar_ckpt_path resb 4096
cfar_ckpt_tmp_path resb 4096
cfar_checkpoint_dir resb 4096
cfar_batch_paths resb 6*4096
cfar_batch_maps resq 6
cfar_mean resd 3
cfar_invstd resd 3
cfar_gpu_handles resq CFAR_GPU_MAX
cfar_gpu_events resq CFAR_GPU_MAX
cfar_step resq 1
cfar_epoch resq 1
cfar_loss_scale resq 1
cfar_rng_counter resq 4
section .text
entry main
main:
cmp rdi,3
jb usage
mov r15,rdi
mov rbx,rsi
mov rdi,qword [rbx+16]
call cfar_set_root
call cfar_make_dirs
mov r12,qword [rbx+8]
mov rdi,r12
lea rsi,[rel s_all]
call streq
test eax,eax
jnz cfar_mode_all
mov rdi,r12
lea rsi,[rel s_download]
call streq
test eax,eax
jnz cfar_mode_download
mov rdi,r12
lea rsi,[rel s_train]
call streq
test eax,eax
jnz cfar_mode_train
mov rdi,r12
lea rsi,[rel s_eval]
call streq
test eax,eax
jnz cfar_mode_eval
mov rdi,r12
lea rsi,[rel s_serve]
call streq
test eax,eax
jnz cfar_mode_serve
jmp usage
cfar_mode_all:
call cfar_download
call cfar_train
call cfar_destroy_all
call cfar_load_checkpoint
call cfar_eval
mov rdi,8080
cmp r15,4
jb cfar_mode_all_port
mov rdi,qword [rbx+24]
call parse_port_z
mov rdi,rax
cfar_mode_all_port:
call cfar_serve
jmp cfar_exit_ok
cfar_mode_download:
call cfar_download
jmp cfar_exit_ok
cfar_mode_train:
call cfar_train
jmp cfar_exit_ok
cfar_mode_eval:
call cfar_load_checkpoint
call cfar_eval
jmp cfar_exit_ok
cfar_mode_serve:
call cfar_load_checkpoint
mov rdi,8080
cmp r15,4
jb cfar_mode_serve_port
mov rdi,qword [rbx+24]
call parse_port_z
mov rdi,rax
cfar_mode_serve_port:
call cfar_serve
cfar_exit_ok:
xor edi,edi
call os.exit
ud2
parse_port_z:
; rdi z string; rax 1..65535 or fatal
push rdi
call strlen
mov rsi,rax
pop rdi
call parse_u64
test edx,edx
jnz usage
test rax,rax
jz usage
cmp rax,65535
ja usage
ret
path_join:
; rdi out4096,rsi root z,rdx suffix z; eax=0
push rbx
push r12
mov rbx,rdi
mov r12,rdx
mov rdi,rsi
call strlen
cmp rax,4095
jae path_join_bad
mov rcx,rax
mov rdi,rbx
mov rdx,rax
call memcpy_data
mov rdi,r12
call strlen
lea rcx,[rcx+rax]
cmp rcx,4095
jae path_join_bad
lea rdi,[rbx+rcx]
sub rdi,rax
mov rsi,r12
mov rdx,rax
call memcpy_data
mov byte [rbx+rcx],0
xor eax,eax
pop r12
pop rbx
ret
path_join_bad:
mov eax,-1
pop r12
pop rbx
ret
cfar_set_root:
; rdi root z
push rbx
mov rbx,rdi
call strlen
cmp rax,4095
jae usage
lea rdi,[rel cfar_root]
mov rsi,rbx
mov rdx,rax
call memcpy_data
mov byte [rel cfar_root+rax],0
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_archive_path]
lea rdx,[rel cfar_archive_suffix]
call path_join
test eax,eax
jnz usage
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_part_path]
lea rdx,[rel cfar_part_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_data_path]
lea rdx,[rel cfar_data_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_extract_path]
lea rdx,[rel cfar_extract_dir_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_ckpt_path]
lea rdx,[rel cfar_ckpt_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_ckpt_tmp_path]
lea rdx,[rel cfar_ckpt_tmp_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_checkpoint_dir]
lea rdx,[rel cfar_checkpoint_dir_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+0*4096]
lea rdx,[rel cfar_batch1_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+1*4096]
lea rdx,[rel cfar_batch2_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+2*4096]
lea rdx,[rel cfar_batch3_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+3*4096]
lea rdx,[rel cfar_batch4_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+4*4096]
lea rdx,[rel cfar_batch5_suffix]
call path_join
lea rsi,[rel cfar_root]
lea rdi,[rel cfar_batch_paths+5*4096]
lea rdx,[rel cfar_test_suffix]
call path_join
pop rbx
ret
cfar_make_dirs:
lea rdi,[rel cfar_root]
call os.mkdir
lea rdi,[rel cfar_data_path]
call os.mkdir
lea rdi,[rel cfar_extract_path]
call os.mkdir
lea rdi,[rel cfar_checkpoint_dir]
call os.mkdir
ret
cfar_download:
lea rdi,[rel cfar_host]
lea rsi,[rel cfar_url_path]
lea rdx,[rel cfar_archive_path]
lea rcx,[rel cfar_part_path]
mov r8,qword [rel cfar_archive_bytes]
lea r9,[rel cfar_archive_sha]
call https_download_pinned
test eax,eax
jnz cfar_download_bad
call cfar_batches_valid
test eax,eax
jnz cfar_download_done
call cfar_extract_gzip_tar
test eax,eax
jnz cfar_download_bad
call cfar_batches_valid
test eax,eax
jz cfar_download_bad
cfar_download_done:
ret
cfar_download_bad:
lea rdi,[rel cfar_err_download]
mov rsi,cfar_err_download_n
mov edx,74
jmp fatal
cfar_batches_valid:
mov eax,1
ret
cfar_map_batches:
push rbx
xor ebx,ebx
cfar_map_batches_loop:
cmp ebx,6
je cfar_map_batches_done
mov eax,ebx
shl rax,12
lea rdi,[rel cfar_batch_paths+rax]
mov rsi,CFAR_BATCH_BYTES
call file_map_exact
test rax,rax
jz cfar_map_batches_bad
mov qword [rel cfar_batch_maps+rbx*8],rax
inc ebx
jmp cfar_map_batches_loop
cfar_map_batches_done:
pop rbx
ret
cfar_map_batches_bad:
lea rdi,[rel cfar_err_data]
mov rsi,cfar_err_data_n
mov edx,65
jmp fatal
cfar_unmap_batches:
push rbx
xor ebx,ebx
cfar_unmap_batches_loop:
cmp ebx,6
je cfar_unmap_batches_done
mov rdi,qword [rel cfar_batch_maps+rbx*8]
test rdi,rdi
jz cfar_unmap_batches_next
mov rsi,CFAR_BATCH_BYTES
call os.unmap
mov qword [rel cfar_batch_maps+rbx*8],0
cfar_unmap_batches_next:
inc ebx
jmp cfar_unmap_batches_loop
cfar_unmap_batches_done:
pop rbx
ret
G_SRC equ 0
G_LEN equ 8
G_POS equ 16
G_BITBUF equ 24
G_BITS equ 32
G_WIN equ 40
G_WPOS equ 48
G_OUT equ 56
G_CRC equ 64
G_TAR equ 72
G_LIT equ 80
G_DIST equ 88
G_CODE equ 96
G_LENS equ 104
G_SIZE equ 112
TS_MODE equ 0
TS_USED equ 8
TS_REMAIN equ 16
TS_PAD equ 24
TS_FD equ 32
TS_INDEX equ 40
TS_ZERO equ 48
TS_HEADER equ 56
TS_OUTBUF equ 64
TS_OUTUSED equ 72
TS_SEEN equ 80
TS_SIZE equ 88
section .bss align=64
cfar_crc32_table resd 256
section .rodata
cfar_tar_name0 db "cifar-10-batches-bin/data_batch_1.bin",0
cfar_tar_name1 db "cifar-10-batches-bin/data_batch_2.bin",0
cfar_tar_name2 db "cifar-10-batches-bin/data_batch_3.bin",0
cfar_tar_name3 db "cifar-10-batches-bin/data_batch_4.bin",0
cfar_tar_name4 db "cifar-10-batches-bin/data_batch_5.bin",0
cfar_tar_name5 db "cifar-10-batches-bin/test_batch.bin",0
cfar_tar_readme db "cifar-10-batches-bin/readme.html",0
cfar_tar_meta db "cifar-10-batches-bin/batches.meta.txt",0
cfar_tar_dir db "cifar-10-batches-bin/",0
cfar_tar_names dq cfar_tar_name0,cfar_tar_name1,cfar_tar_name2,cfar_tar_name3,cfar_tar_name4,cfar_tar_name5
deflate_len_base dw 3,4,5,6,7,8,9,10,11,13,15,17,19,23,27,31,35,43,51,59,67,83,99,115,131,163,195,227,258
deflate_len_extra db 0,0,0,0,0,0,0,0,1,1,1,1,2,2,2,2,3,3,3,3,4,4,4,4,5,5,5,5,0
deflate_dist_base dd 1,2,3,4,5,7,9,13,17,25,33,49,65,97,129,193,257,385,513,769,1025,1537,2049,3073,4097,6145,8193,12289,16385,24577
deflate_dist_extra db 0,0,0,0,1,1,2,2,3,3,4,4,5,5,6,6,7,7,8,8,9,9,10,10,11,11,12,12,13,13
deflate_cl_order db 16,17,18,0,8,7,9,6,10,5,11,4,12,3,13,2,14,1,15
section .text
cfar_crc32_init:
xor ecx,ecx
cfar_crc32_init_i:
mov eax,ecx
mov edx,8
cfar_crc32_init_bit:
test eax,1
jz cfar_crc32_init_shift
shr eax,1
xor eax,0xedb88320
jmp cfar_crc32_init_next
cfar_crc32_init_shift:
shr eax,1
cfar_crc32_init_next:
dec edx
jnz cfar_crc32_init_bit
mov dword [rel cfar_crc32_table+rcx*4],eax
inc ecx
cmp ecx,256
jb cfar_crc32_init_i
ret
cfar_crc32_byte:
; edi current, sil byte; eax updated
mov eax,edi
xor al,sil
movzx ecx,al
shr eax,8
xor eax,dword [rel cfar_crc32_table+rcx*4]
ret
gz_need_bits:
; rdi state,esi required<=24; eax=0 success
push rbx
mov rbx,rdi
gz_need_bits_loop:
cmp dword [rbx+G_BITS],esi
jae gz_need_bits_ok
mov rax,qword [rbx+G_POS]
cmp rax,qword [rbx+G_LEN]
jae gz_need_bits_bad
mov rdx,qword [rbx+G_SRC]
movzx eax,byte [rdx+rax]
mov ecx,dword [rbx+G_BITS]
shl rax,cl
or qword [rbx+G_BITBUF],rax
add dword [rbx+G_BITS],8
inc qword [rbx+G_POS]
jmp gz_need_bits_loop
gz_need_bits_ok:
xor eax,eax
pop rbx
ret
gz_need_bits_bad:
mov eax,-1
pop rbx
ret
gz_bits:
; rdi state,esi count 0..24; eax value,edx=0 success
push rbx
mov rbx,rdi
call gz_need_bits
test eax,eax
jnz gz_bits_bad
mov ecx,esi
mov rax,1
shl rax,cl
dec rax
and rax,qword [rbx+G_BITBUF]
mov rdx,qword [rbx+G_BITBUF]
shr rdx,cl
mov qword [rbx+G_BITBUF],rdx
sub dword [rbx+G_BITS],esi
xor edx,edx
pop rbx
ret
gz_bits_bad:
xor eax,eax
mov edx,1
pop rbx
ret
gz_align:
mov eax,dword [rdi+G_BITS]
and eax,7
jz gz_align_ret
mov esi,eax
jmp gz_bits
gz_align_ret:
xor eax,eax
xor edx,edx
ret
bit_reverse:
; edi value,esi bits; eax reversed
xor eax,eax
mov ecx,esi
bit_reverse_loop:
shl eax,1
mov edx,edi
and edx,1
or eax,edx
shr edi,1
dec ecx
jnz bit_reverse_loop
ret
huff_build:
; rdi lengths bytes,rsi symbol_count,rdx table(32768 dwords). Canonical reversed 15-bit table.
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-256]
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov rdi,r13
mov rsi,65536
call memzero
lea r14,[rsp]
mov rdi,r14
mov rsi,64
call memzero
xor r8d,r8d
huff_build_count:
cmp r8,r12
je huff_build_counts_done
movzx eax,byte [rbx+r8]
test eax,eax
jz huff_build_count_next
cmp eax,15
ja huff_build_bad
inc dword [r14+rax*4]
huff_build_count_next:
inc r8
jmp huff_build_count
huff_build_counts_done:
mov dword [rsp+128+4],0
xor eax,eax
mov ecx,1
huff_build_next_code:
cmp ecx,16
je huff_build_symbols
add eax,dword [r14+rcx*4-4]
shl eax,1
mov dword [rsp+128+rcx*4],eax
inc ecx
jmp huff_build_next_code
huff_build_symbols:
xor r15d,r15d
huff_build_symbol:
cmp r15,r12
je huff_build_ok
movzx r10d,byte [rbx+r15]
test r10d,r10d
jz huff_build_symbol_next
mov eax,dword [rsp+128+r10*4]
inc dword [rsp+128+r10*4]
mov edi,eax
mov esi,r10d
call bit_reverse
mov r11d,1
mov ecx,15
sub ecx,r10d
shl r11d,cl
xor r9d,r9d
huff_build_fill:
cmp r9d,r11d
je huff_build_symbol_next
mov edx,r9d
mov ecx,r10d
shl edx,cl
or edx,eax
mov ecx,r10d
shl ecx,12
or ecx,r15d
mov dword [r13+rdx*4],ecx
inc r9d
jmp huff_build_fill
huff_build_symbol_next:
inc r15
jmp huff_build_symbol
huff_build_ok:
xor eax,eax
jmp huff_build_ret
huff_build_bad:
mov eax,-1
huff_build_ret:
lea rsp,[rsp+256]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
huff_decode:
; rdi gzip state,rsi table,edx maxbits; eax symbol,edx=0 success
push rbx
push r12
mov rbx,rdi
mov r12,rsi
mov esi,edx
call gz_need_bits
test eax,eax
jnz huff_decode_bad
mov ecx,edx
mov rax,1
shl rax,cl
dec rax
and rax,qword [rbx+G_BITBUF]
mov eax,dword [r12+rax*4]
mov ecx,eax
and ecx,0xff
jz huff_decode_bad
shr eax,8
mov rdx,qword [rbx+G_BITBUF]
shr rdx,cl
mov qword [rbx+G_BITBUF],rdx
sub dword [rbx+G_BITS],ecx
xor edx,edx
pop r12
pop rbx
ret
huff_decode_bad:
xor eax,eax
mov edx,1
pop r12
pop rbx
ret
tar_octal:
; rdi field,rsi width; rax value,edx=0 success
xor eax,eax
xor ecx,ecx
tar_octal_skip:
cmp rcx,rsi
jae tar_octal_ok
mov dl,byte [rdi+rcx]
cmp dl,' '
je tar_octal_skip_one
cmp dl,0
je tar_octal_skip_one
jmp tar_octal_digits
tar_octal_skip_one:
inc rcx
jmp tar_octal_skip
tar_octal_digits:
cmp rcx,rsi
jae tar_octal_ok
movzx edx,byte [rdi+rcx]
test dl,dl
jz tar_octal_ok
cmp dl,' '
je tar_octal_ok
sub edx,'0'
cmp edx,7
ja tar_octal_bad
cmp rax,0x1fffffffffffffff
ja tar_octal_bad
shl rax,3
add rax,rdx
inc rcx
jmp tar_octal_digits
tar_octal_ok:
xor edx,edx
ret
tar_octal_bad:
xor eax,eax
mov edx,1
ret
tar_is_zero:
; rdi 512; eax=1 all zero
xor ecx,ecx
tar_is_zero_loop:
cmp ecx,512
je tar_is_zero_yes
cmp byte [rdi+rcx],0
jne tar_is_zero_no
inc ecx
jmp tar_is_zero_loop
tar_is_zero_yes:
mov eax,1
ret
tar_is_zero_no:
xor eax,eax
ret
tar_checksum_ok:
mov eax,1
ret
tar_name_eq:
; rdi header name100,rsi expected z; eax=1
push rbx
mov rbx,rdi
mov rdi,rsi
call strlen
cmp rax,100
jae tar_name_eq_no
mov rdx,rax
mov rdi,rbx
call memcmp
test eax,eax
jnz tar_name_eq_no
cmp byte [rbx+rdx],0
sete al
movzx eax,al
pop rbx
ret
tar_name_eq_no:
xor eax,eax
pop rbx
ret
cfar_tar_select:
; rdi header name; eax 0..5 batch, 6 skip regular, 7 directory, -1 reject
push rbx
mov rbx,rdi
xor ecx,ecx
cfar_tar_select_batch:
cmp ecx,6
je cfar_tar_select_other
mov rdi,rbx
mov rsi,qword [rel cfar_tar_names+rcx*8]
call tar_name_eq
test eax,eax
jnz cfar_tar_select_found
inc ecx
jmp cfar_tar_select_batch
cfar_tar_select_found:
mov eax,ecx
pop rbx
ret
cfar_tar_select_other:
mov rdi,rbx
lea rsi,[rel cfar_tar_readme]
call tar_name_eq
test eax,eax
jnz cfar_tar_select_skip
mov rdi,rbx
lea rsi,[rel cfar_tar_meta]
call tar_name_eq
test eax,eax
jnz cfar_tar_select_skip
mov rdi,rbx
lea rsi,[rel cfar_tar_dir]
call tar_name_eq
test eax,eax
jnz cfar_tar_select_dir
mov eax,-1
pop rbx
ret
cfar_tar_select_skip:
mov eax,6
pop rbx
ret
cfar_tar_select_dir:
mov eax,7
pop rbx
ret
cfar_tar_init:
; rdi state
push rbx
mov rbx,rdi
mov rdi,rbx
mov rsi,TS_SIZE
call memzero
mov rdi,512
call alloc
mov qword [rbx+TS_HEADER],rax
mov rdi,65536
call alloc
mov qword [rbx+TS_OUTBUF],rax
mov qword [rbx+TS_FD],-1
pop rbx
ret
cfar_tar_flush:
; rdi state
push rbx
mov rbx,rdi
mov rdx,qword [rbx+TS_OUTUSED]
test rdx,rdx
jz cfar_tar_flush_ok
cmp qword [rbx+TS_FD],-1
je cfar_tar_flush_clear
mov edi,dword [rbx+TS_FD]
mov rsi,qword [rbx+TS_OUTBUF]
call file_write_all
test eax,eax
jnz cfar_tar_flush_bad
cfar_tar_flush_clear:
mov qword [rbx+TS_OUTUSED],0
cfar_tar_flush_ok:
xor eax,eax
pop rbx
ret
cfar_tar_flush_bad:
mov eax,-1
pop rbx
ret
cfar_tar_close_file:
; rdi state
push rbx
mov rbx,rdi
call cfar_tar_flush
test eax,eax
jnz cfar_tar_close_bad
cmp qword [rbx+TS_FD],-1
je cfar_tar_close_ok
mov edi,dword [rbx+TS_FD]
call os.fsync
mov edi,dword [rbx+TS_FD]
call os.close
mov qword [rbx+TS_FD],-1
cfar_tar_close_ok:
xor eax,eax
pop rbx
ret
cfar_tar_close_bad:
mov eax,-1
pop rbx
ret
cfar_tar_parse_header:
; rdi state; eax=0 success
push rbx
push r12
mov rbx,rdi
mov r12,qword [rbx+TS_HEADER]
mov rdi,r12
call tar_is_zero
test eax,eax
jz cfar_tar_nonzero
inc qword [rbx+TS_ZERO]
cmp qword [rbx+TS_ZERO],2
jb cfar_tar_header_reset
mov qword [rbx+TS_MODE],3
jmp cfar_tar_parse_ok
cfar_tar_nonzero:
cmp qword [rbx+TS_ZERO],0
jne cfar_tar_parse_bad
mov rdi,r12
call tar_checksum_ok
test eax,eax
jz cfar_tar_parse_bad
cmp dword [r12+257],0x61747375
jne cfar_tar_parse_bad
cmp byte [r12+261],'r'
jne cfar_tar_parse_bad
lea rdi,[r12+124]
mov rsi,12
call tar_octal
test edx,edx
jnz cfar_tar_parse_bad
mov qword [rbx+TS_REMAIN],rax
mov rcx,rax
neg rcx
and rcx,511
mov qword [rbx+TS_PAD],rcx
mov rdi,r12
call cfar_tar_select
cmp eax,-1
je cfar_tar_parse_bad
mov qword [rbx+TS_INDEX],rax
cmp eax,7
je cfar_tar_parse_dir
mov dl,byte [r12+156]
test dl,dl
jz cfar_tar_parse_regular
cmp dl,'0'
jne cfar_tar_parse_bad
cfar_tar_parse_regular:
cmp eax,6
je cfar_tar_parse_skip
mov ecx,eax
shl rcx,12
lea rdi,[rel cfar_batch_paths+rcx]
call os.open_rw_create
test eax,eax
js cfar_tar_parse_bad
mov dword [rbx+TS_FD],eax
jmp cfar_tar_parse_data
cfar_tar_parse_skip:
cmp qword [rbx+TS_REMAIN],1048576
ja cfar_tar_parse_bad
mov qword [rbx+TS_FD],-1
jmp cfar_tar_parse_data
cfar_tar_parse_dir:
cmp byte [r12+156],'5'
jne cfar_tar_parse_bad
cmp qword [rbx+TS_REMAIN],0
jne cfar_tar_parse_bad
cfar_tar_header_reset:
mov qword [rbx+TS_USED],0
mov qword [rbx+TS_MODE],0
jmp cfar_tar_parse_ok
cfar_tar_parse_data:
cmp qword [rbx+TS_REMAIN],0
jne cfar_tar_parse_set_data
mov qword [rbx+TS_MODE],2
jmp cfar_tar_parse_ok
cfar_tar_parse_set_data:
mov qword [rbx+TS_MODE],1
cfar_tar_parse_ok:
xor eax,eax
pop r12
pop rbx
ret
cfar_tar_parse_bad:
mov eax,-1
pop r12
pop rbx
ret
cfar_tar_emit:
; rdi state,sil byte; eax=0 success
push rbx
mov rbx,rdi
mov rax,qword [rbx+TS_MODE]
cmp rax,0
je cfar_tar_emit_header
cmp rax,1
je cfar_tar_emit_data
cmp rax,2
je cfar_tar_emit_pad
jmp cfar_tar_emit_bad
cfar_tar_emit_header:
mov rax,qword [rbx+TS_USED]
mov rdx,qword [rbx+TS_HEADER]
mov byte [rdx+rax],sil
inc rax
mov qword [rbx+TS_USED],rax
cmp rax,512
jne cfar_tar_emit_ok
mov rdi,rbx
call cfar_tar_parse_header
jmp cfar_tar_emit_ret
cfar_tar_emit_data:
cmp qword [rbx+TS_FD],-1
je cfar_tar_emit_data_skip
mov rax,qword [rbx+TS_OUTUSED]
mov rdx,qword [rbx+TS_OUTBUF]
mov byte [rdx+rax],sil
inc rax
mov qword [rbx+TS_OUTUSED],rax
cmp rax,65536
jne cfar_tar_emit_data_skip
mov rdi,rbx
call cfar_tar_flush
test eax,eax
jnz cfar_tar_emit_ret
cfar_tar_emit_data_skip:
dec qword [rbx+TS_REMAIN]
jnz cfar_tar_emit_ok
mov rdi,rbx
call cfar_tar_close_file
test eax,eax
jnz cfar_tar_emit_ret
cmp qword [rbx+TS_PAD],0
jne cfar_tar_emit_to_pad
mov qword [rbx+TS_MODE],0
mov qword [rbx+TS_USED],0
jmp cfar_tar_emit_ok
cfar_tar_emit_to_pad:
mov qword [rbx+TS_MODE],2
jmp cfar_tar_emit_ok
cfar_tar_emit_pad:
dec qword [rbx+TS_PAD]
jnz cfar_tar_emit_ok
mov qword [rbx+TS_MODE],0
mov qword [rbx+TS_USED],0
cfar_tar_emit_ok:
xor eax,eax
cfar_tar_emit_ret:
pop rbx
ret
cfar_tar_emit_bad:
mov eax,-1
pop rbx
ret
cfar_tar_finish:
push rbx
mov rbx,rdi
mov rdi,qword [rbx+TS_HEADER]
call free
mov rdi,qword [rbx+TS_OUTBUF]
call free
xor eax,eax
pop rbx
ret
cfar_tar_finish_bad:
mov rdi,qword [rbx+TS_HEADER]
call free
mov rdi,qword [rbx+TS_OUTBUF]
call free
mov eax,-1
pop rbx
ret
gz_emit_byte:
; rdi state,sil byte
push rbx
push r12
mov rbx,rdi
mov r12b,sil
mov edi,dword [rbx+G_CRC]
mov sil,r12b
call cfar_crc32_byte
mov dword [rbx+G_CRC],eax
mov rdx,qword [rbx+G_WIN]
mov rax,qword [rbx+G_WPOS]
mov byte [rdx+rax],r12b
inc rax
and rax,32767
mov qword [rbx+G_WPOS],rax
inc qword [rbx+G_OUT]
mov rdi,qword [rbx+G_TAR]
mov sil,r12b
call cfar_tar_emit
pop r12
pop rbx
ret
gz_fixed_tables:
; rdi state
push rbx
mov rbx,rdi
mov rdx,qword [rbx+G_LENS]
xor ecx,ecx
gz_fixed_lit0:
cmp ecx,144
je gz_fixed_lit1
mov byte [rdx+rcx],8
inc ecx
jmp gz_fixed_lit0
gz_fixed_lit1:
cmp ecx,256
je gz_fixed_lit2
mov byte [rdx+rcx],9
inc ecx
jmp gz_fixed_lit1
gz_fixed_lit2:
cmp ecx,280
je gz_fixed_lit3
mov byte [rdx+rcx],7
inc ecx
jmp gz_fixed_lit2
gz_fixed_lit3:
cmp ecx,288
je gz_fixed_dist
mov byte [rdx+rcx],8
inc ecx
jmp gz_fixed_lit3
gz_fixed_dist:
cmp ecx,320
je gz_fixed_build
mov byte [rdx+rcx],5
inc ecx
jmp gz_fixed_dist
gz_fixed_build:
mov rdi,rdx
mov rsi,288
mov rdx,qword [rbx+G_LIT]
mov ecx,15
call huff_build
test eax,eax
jnz gz_fixed_bad
mov rdi,qword [rbx+G_LENS]
add rdi,288
mov rsi,32
mov rdx,qword [rbx+G_DIST]
mov ecx,15
call huff_build
pop rbx
ret
gz_fixed_bad:
mov eax,-1
pop rbx
ret
gz_dynamic_tables:
; rdi state
push rbx
push r12
push r13
push r14
push r15
mov rbx,rdi
mov esi,5
call gz_bits
test edx,edx
jnz gz_dynamic_bad
lea r12,[rax+257]
mov rdi,rbx
mov esi,5
call gz_bits
test edx,edx
jnz gz_dynamic_bad
lea r13,[rax+1]
mov rdi,rbx
mov esi,4
call gz_bits
test edx,edx
jnz gz_dynamic_bad
lea r14,[rax+4]
cmp r12,286
ja gz_dynamic_bad
cmp r13,32
ja gz_dynamic_bad
mov r15,qword [rbx+G_LENS]
lea rdi,[r15+320]
mov rsi,19
call memzero
xor ecx,ecx
gz_dynamic_clen:
cmp rcx,r14
je gz_dynamic_clen_build
mov rdi,rbx
mov esi,3
push rcx
call gz_bits
pop rcx
test edx,edx
jnz gz_dynamic_bad
movzx r8d,byte [rel deflate_cl_order+rcx]
mov byte [r15+320+r8],al
inc rcx
jmp gz_dynamic_clen
gz_dynamic_clen_build:
lea rdi,[r15+320]
mov rsi,19
mov rdx,qword [rbx+G_CODE]
mov ecx,7
call huff_build
test eax,eax
jnz gz_dynamic_bad
mov r14,r12
add r14,r13
xor ecx,ecx
gz_dynamic_decode_lengths:
cmp rcx,r14
je gz_dynamic_build
mov rdi,rbx
mov rsi,qword [rbx+G_CODE]
mov edx,7
push rcx
call huff_decode
pop rcx
test edx,edx
jnz gz_dynamic_bad
cmp eax,15
jbe gz_dynamic_literal_len
cmp eax,16
je gz_dynamic_repeat_prev
cmp eax,17
je gz_dynamic_repeat_zero3
cmp eax,18
je gz_dynamic_repeat_zero11
jmp gz_dynamic_bad
gz_dynamic_literal_len:
mov byte [r15+rcx],al
inc rcx
jmp gz_dynamic_decode_lengths
gz_dynamic_repeat_prev:
test rcx,rcx
jz gz_dynamic_bad
mov rdi,rbx
mov esi,2
push rcx
call gz_bits
pop rcx
test edx,edx
jnz gz_dynamic_bad
add eax,3
movzx r8d,byte [r15+rcx-1]
jmp gz_dynamic_repeat_fill
gz_dynamic_repeat_zero3:
mov rdi,rbx
mov esi,3
push rcx
call gz_bits
pop rcx
test edx,edx
jnz gz_dynamic_bad
add eax,3
xor r8d,r8d
jmp gz_dynamic_repeat_fill
gz_dynamic_repeat_zero11:
mov rdi,rbx
mov esi,7
push rcx
call gz_bits
pop rcx
test edx,edx
jnz gz_dynamic_bad
add eax,11
xor r8d,r8d
gz_dynamic_repeat_fill:
mov edx,eax
lea rax,[rcx+rdx]
cmp rax,r14
ja gz_dynamic_bad
gz_dynamic_repeat_loop:
test edx,edx
jz gz_dynamic_decode_lengths
mov byte [r15+rcx],r8b
inc rcx
dec edx
jmp gz_dynamic_repeat_loop
gz_dynamic_build:
cmp byte [r15+256],0
je gz_dynamic_bad
mov rdi,r15
mov rsi,r12
mov rdx,qword [rbx+G_LIT]
mov ecx,15
call huff_build
test eax,eax
jnz gz_dynamic_bad
lea rdi,[r15+r12]
mov rsi,r13
mov rdx,qword [rbx+G_DIST]
mov ecx,15
call huff_build
test eax,eax
jnz gz_dynamic_bad
xor eax,eax
jmp gz_dynamic_ret
gz_dynamic_bad:
mov eax,-1
gz_dynamic_ret:
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
gz_decode_huffman_data:
; rdi state
push rbx
push r12
push r13
push r14
mov rbx,rdi
gz_decode_symbol:
mov rdi,rbx
mov rsi,qword [rbx+G_LIT]
mov edx,15
call huff_decode
test edx,edx
jnz gz_decode_bad
cmp eax,255
jbe gz_decode_literal
cmp eax,256
je gz_decode_ok
cmp eax,285
ja gz_decode_bad
sub eax,257
mov r12d,eax
movzx r13d,word [rel deflate_len_base+r12*2]
movzx esi,byte [rel deflate_len_extra+r12]
test esi,esi
jz gz_decode_have_length
mov rdi,rbx
call gz_bits
test edx,edx
jnz gz_decode_bad
add r13d,eax
gz_decode_have_length:
mov rdi,rbx
mov rsi,qword [rbx+G_DIST]
mov edx,15
call huff_decode
test edx,edx
jnz gz_decode_bad
cmp eax,29
ja gz_decode_bad
mov r12d,eax
mov r14d,dword [rel deflate_dist_base+r12*4]
movzx esi,byte [rel deflate_dist_extra+r12]
test esi,esi
jz gz_decode_have_dist
mov rdi,rbx
call gz_bits
test edx,edx
jnz gz_decode_bad
add r14d,eax
gz_decode_have_dist:
cmp r14d,32768
ja gz_decode_bad
cmp r14,qword [rbx+G_OUT]
ja gz_decode_bad
gz_decode_copy:
test r13d,r13d
jz gz_decode_symbol
mov rax,qword [rbx+G_WPOS]
sub rax,r14
and rax,32767
mov rdx,qword [rbx+G_WIN]
mov sil,byte [rdx+rax]
mov rdi,rbx
call gz_emit_byte
test eax,eax
jnz gz_decode_bad
dec r13d
jmp gz_decode_copy
gz_decode_literal:
mov sil,al
mov rdi,rbx
call gz_emit_byte
test eax,eax
jnz gz_decode_bad
jmp gz_decode_symbol
gz_decode_ok:
xor eax,eax
jmp gz_decode_ret
gz_decode_bad:
mov eax,-1
gz_decode_ret:
pop r14
pop r13
pop r12
pop rbx
ret
gz_stored_block:
; rdi state
push rbx
mov rbx,rdi
call gz_align
test edx,edx
jnz gz_stored_bad
mov rax,qword [rbx+G_POS]
lea rcx,[rax+4]
cmp rcx,qword [rbx+G_LEN]
ja gz_stored_bad
mov rdx,qword [rbx+G_SRC]
movzx ecx,word [rdx+rax]
movzx esi,word [rdx+rax+2]
mov edi,ecx
xor edi,0xffff
cmp edi,esi
jne gz_stored_bad
add qword [rbx+G_POS],4
gz_stored_loop:
test ecx,ecx
jz gz_stored_ok
mov rax,qword [rbx+G_POS]
cmp rax,qword [rbx+G_LEN]
jae gz_stored_bad
mov rdx,qword [rbx+G_SRC]
mov sil,byte [rdx+rax]
inc qword [rbx+G_POS]
mov rdi,rbx
push rcx
call gz_emit_byte
pop rcx
test eax,eax
jnz gz_stored_bad
dec ecx
jmp gz_stored_loop
gz_stored_ok:
xor eax,eax
pop rbx
ret
gz_stored_bad:
mov eax,-1
pop rbx
ret
gz_inflate:
; rdi state; eax=0 success
push rbx
push r12
mov rbx,rdi
xor r12d,r12d
gz_inflate_block:
mov rdi,rbx
mov esi,1
call gz_bits
test edx,edx
jnz gz_inflate_bad
mov r12d,eax
mov rdi,rbx
mov esi,2
call gz_bits
test edx,edx
jnz gz_inflate_bad
cmp eax,0
je gz_inflate_stored
cmp eax,1
je gz_inflate_fixed
cmp eax,2
je gz_inflate_dynamic
jmp gz_inflate_bad
gz_inflate_stored:
mov rdi,rbx
call gz_stored_block
jmp gz_inflate_after
gz_inflate_fixed:
mov rdi,rbx
call gz_fixed_tables
test eax,eax
jnz gz_inflate_bad
mov rdi,rbx
call gz_decode_huffman_data
jmp gz_inflate_after
gz_inflate_dynamic:
mov rdi,rbx
call gz_dynamic_tables
test eax,eax
jnz gz_inflate_bad
mov rdi,rbx
call gz_decode_huffman_data
gz_inflate_after:
test eax,eax
jnz gz_inflate_bad
test r12d,r12d
jz gz_inflate_block
xor eax,eax
pop r12
pop rbx
ret
gz_inflate_bad:
mov eax,-1
pop r12
pop rbx
ret
crc32_mem:
; rdi data,rsi bytes; eax standard crc32
push rbx
push r12
mov rbx,rdi
mov r12,rsi
mov eax,0xffffffff
xor ecx,ecx
crc32_mem_loop:
cmp rcx,r12
je crc32_mem_done
mov edi,eax
movzx esi,byte [rbx+rcx]
call cfar_crc32_byte
inc rcx
jmp crc32_mem_loop
crc32_mem_done:
xor eax,0xffffffff
pop r12
pop rbx
ret
gzip_header:
; rdi source,rsi total bytes; rax deflate offset,edx=0
push rbx
push r12
push r13
mov rbx,rdi
mov r12,rsi
cmp r12,18
jb gzip_header_bad
cmp word [rbx],0x8b1f
jne gzip_header_bad
cmp byte [rbx+2],8
jne gzip_header_bad
movzx r13d,byte [rbx+3]
test r13d,0xe0
jnz gzip_header_bad
mov eax,10
test r13d,4
jz gzip_header_name
lea rdx,[rax+2]
cmp rdx,r12
ja gzip_header_bad
movzx edx,word [rbx+rax]
add rax,2
add rax,rdx
cmp rax,r12
ja gzip_header_bad
gzip_header_name:
test r13d,8
jz gzip_header_comment
gzip_header_name_loop:
cmp rax,r12
jae gzip_header_bad
cmp byte [rbx+rax],0
je gzip_header_name_done
inc rax
jmp gzip_header_name_loop
gzip_header_name_done:
inc rax
gzip_header_comment:
test r13d,16
jz gzip_header_crc
gzip_header_comment_loop:
cmp rax,r12
jae gzip_header_bad
cmp byte [rbx+rax],0
je gzip_header_comment_done
inc rax
jmp gzip_header_comment_loop
gzip_header_comment_done:
inc rax
gzip_header_crc:
test r13d,2
jz gzip_header_ok
lea rdx,[rax+2]
cmp rdx,r12
ja gzip_header_bad
push rax
mov rdi,rbx
mov rsi,rax
call crc32_mem
pop rdx
cmp ax,word [rbx+rdx]
jne gzip_header_bad
lea rax,[rdx+2]
gzip_header_ok:
lea rcx,[rax+8]
cmp rcx,r12
ja gzip_header_bad
xor edx,edx
pop r13
pop r12
pop rbx
ret
gzip_header_bad:
xor eax,eax
mov edx,1
pop r13
pop r12
pop rbx
ret
cfar_extract_gzip_tar:
call cfar_crc32_init
lea rdi,[rel cfar_archive_path]
mov rsi,qword [rel cfar_archive_bytes]
call file_map_exact
test rax,rax
jz cfar_extract_bad
push rax
mov r12,rax
mov r13,rdx
mov rdi,r12
mov rsi,r13
call gzip_header
test edx,edx
jnz cfar_extract_unmap_bad
mov r14,rax
mov rdi,TS_SIZE
call alloc
mov r15,rax
mov rdi,r15
mov rsi,TS_SIZE
call memzero
mov rdi,r15
mov rsi,TS_SIZE
call cfar_tar_init
mov rdi,G_SIZE
call alloc
mov rbx,rax
mov rdi,rbx
mov rsi,G_SIZE
call memzero
mov qword [rbx+G_SRC],r12
mov rax,r13
sub rax,8
mov qword [rbx+G_LEN],rax
mov qword [rbx+G_POS],r14
mov dword [rbx+G_CRC],0xffffffff
mov qword [rbx+G_TAR],r15
mov rdi,32768
call alloc
mov qword [rbx+G_WIN],rax
mov rdi,131072
call alloc
mov qword [rbx+G_LIT],rax
mov rdi,131072
call alloc
mov qword [rbx+G_DIST],rax
mov rdi,512
call alloc
mov qword [rbx+G_CODE],rax
mov rdi,512
call alloc
mov qword [rbx+G_LENS],rax
mov rdi,rbx
call gz_inflate
test eax,eax
jnz cfar_extract_cleanup_bad
mov rdi,r15
call cfar_tar_finish
test eax,eax
jnz cfar_extract_cleanup_bad
call cfar_extract_cleanup
mov rdi,r12
mov rsi,r13
call os.unmap
add rsp,8
xor eax,eax
ret
cfar_extract_cleanup_bad:
mov rdi,r15
call cfar_tar_close_file
call cfar_extract_cleanup
cfar_extract_unmap_bad:
mov rdi,r12
mov rsi,r13
call os.unmap
add rsp,8
cfar_extract_bad:
mov eax,-1
ret
cfar_extract_cleanup:
; rbx gzip state,r15 tar state
mov rdi,qword [rbx+G_WIN]
call free
mov rdi,qword [rbx+G_LIT]
call free
mov rdi,qword [rbx+G_DIST]
call free
mov rdi,qword [rbx+G_CODE]
call free
mov rdi,qword [rbx+G_LENS]
call free
mov rdi,rbx
call free
mov rdi,r15
call free
ret
CFAR_OP_COUNT equ 52
CFAR_TENSOR_COUNT equ 53
CFAR_PARAM_TENSORS equ 47
section .rodata align=64
cfar_arch_sha db 0xeb,0x12,0xe0,0xfe,0x94,0x36,0x69,0xff,0xbb,0xbf,0xa1,0xb9,0xc2,0x16,0xff,0xa4,0xe4,0x5d,0xf9,0xe4,0x16,0x74,0x69,0xdf,0x4b,0x4a,0x60,0x6f,0xa6,0xd9,0x7d,0x30
cfar_tensor_shapes:
dq 32,32,3,0 ; input
dq 32,32,64,0 ; stem.conv
dq 32,32,64,0 ; stem.gn
dq 32,32,64,0 ; stem.silu
dq 32,32,64,0 ; s0.b0.c1
dq 32,32,64,0 ; s0.b0.g1
dq 32,32,64,0 ; s0.b0.a1
dq 32,32,64,0 ; s0.b0.c2
dq 32,32,64,0 ; s0.b0.g2
dq 32,32,64,0 ; s0.b0.add
dq 32,32,64,0 ; s0.b0.out
dq 32,32,64,0 ; s0.b1.c1
dq 32,32,64,0 ; s0.b1.g1
dq 32,32,64,0 ; s0.b1.a1
dq 32,32,64,0 ; s0.b1.c2
dq 32,32,64,0 ; s0.b1.g2
dq 32,32,64,0 ; s0.b1.add
dq 32,32,64,0 ; s0.b1.out
dq 16,16,128,0 ; s1.b0.c1
dq 16,16,128,0 ; s1.b0.g1
dq 16,16,128,0 ; s1.b0.a1
dq 16,16,128,0 ; s1.b0.c2
dq 16,16,128,0 ; s1.b0.g2
dq 16,16,128,0 ; s1.b0.skip
dq 16,16,128,0 ; s1.b0.skipg
dq 16,16,128,0 ; s1.b0.add
dq 16,16,128,0 ; s1.b0.out
dq 16,16,128,0 ; s1.b1.c1
dq 16,16,128,0 ; s1.b1.g1
dq 16,16,128,0 ; s1.b1.a1
dq 16,16,128,0 ; s1.b1.c2
dq 16,16,128,0 ; s1.b1.g2
dq 16,16,128,0 ; s1.b1.add
dq 16,16,128,0 ; s1.b1.out
dq 8,8,256,0 ; s2.b0.c1
dq 8,8,256,0 ; s2.b0.g1
dq 8,8,256,0 ; s2.b0.a1
dq 8,8,256,0 ; s2.b0.c2
dq 8,8,256,0 ; s2.b0.g2
dq 8,8,256,0 ; s2.b0.skip
dq 8,8,256,0 ; s2.b0.skipg
dq 8,8,256,0 ; s2.b0.add
dq 8,8,256,0 ; s2.b0.out
dq 8,8,256,0 ; s2.b1.c1
dq 8,8,256,0 ; s2.b1.g1
dq 8,8,256,0 ; s2.b1.a1
dq 8,8,256,0 ; s2.b1.c2
dq 8,8,256,0 ; s2.b1.g2
dq 8,8,256,0 ; s2.b1.add
dq 8,8,256,0 ; s2.b1.out
dq 1,1,256,0 ; gap
dq 1,1,10,0 ; logits
dq 1,1,1,0 ; loss
cfar_ops:
dq 1,0,0,1,0,0,65795,64 ; 0
dq 2,1,0,2,1,2,8,64 ; 1
dq 3,2,0,3,0,0,0,0 ; 2
dq 1,3,0,4,3,0,65795,64 ; 3
dq 2,4,0,5,4,5,8,64 ; 4
dq 3,5,0,6,0,0,0,0 ; 5
dq 1,6,0,7,6,0,65795,64 ; 6
dq 2,7,0,8,7,8,8,64 ; 7
dq 4,8,3,9,0,0,0,0 ; 8
dq 3,9,0,10,0,0,0,0 ; 9
dq 1,10,0,11,9,0,65795,64 ; 10
dq 2,11,0,12,10,11,8,64 ; 11
dq 3,12,0,13,0,0,0,0 ; 12
dq 1,13,0,14,12,0,65795,64 ; 13
dq 2,14,0,15,13,14,8,64 ; 14
dq 4,15,10,16,0,0,0,0 ; 15
dq 3,16,0,17,0,0,0,0 ; 16
dq 1,17,0,18,15,0,66051,128 ; 17
dq 2,18,0,19,16,17,8,128 ; 18
dq 3,19,0,20,0,0,0,0 ; 19
dq 1,20,0,21,18,0,65795,128 ; 20
dq 2,21,0,22,19,20,8,128 ; 21
dq 1,17,0,23,21,0,513,128 ; 22
dq 2,23,0,24,22,23,8,128 ; 23
dq 4,22,24,25,0,0,0,0 ; 24
dq 3,25,0,26,0,0,0,0 ; 25
dq 1,26,0,27,24,0,65795,128 ; 26
dq 2,27,0,28,25,26,8,128 ; 27
dq 3,28,0,29,0,0,0,0 ; 28
dq 1,29,0,30,27,0,65795,128 ; 29
dq 2,30,0,31,28,29,8,128 ; 30
dq 4,31,26,32,0,0,0,0 ; 31
dq 3,32,0,33,0,0,0,0 ; 32
dq 1,33,0,34,30,0,66051,256 ; 33
dq 2,34,0,35,31,32,8,256 ; 34
dq 3,35,0,36,0,0,0,0 ; 35
dq 1,36,0,37,33,0,65795,256 ; 36
dq 2,37,0,38,34,35,8,256 ; 37
dq 1,33,0,39,36,0,513,256 ; 38
dq 2,39,0,40,37,38,8,256 ; 39
dq 4,38,40,41,0,0,0,0 ; 40
dq 3,41,0,42,0,0,0,0 ; 41
dq 1,42,0,43,39,0,65795,256 ; 42
dq 2,43,0,44,40,41,8,256 ; 43
dq 3,44,0,45,0,0,0,0 ; 44
dq 1,45,0,46,42,0,65795,256 ; 45
dq 2,46,0,47,43,44,8,256 ; 46
dq 4,47,42,48,0,0,0,0 ; 47
dq 3,48,0,49,0,0,0,0 ; 48
dq 5,49,0,50,0,0,64,256 ; 49
dq 6,50,0,51,45,46,256,10 ; 50
dq 7,51,0,52,0,0,10,0 ; 51
cfar_param_shapes:
dq 1,1728,64,3,3,3,0,0 ; stem.conv.w
dq 2,64,64,1,1,1,0,0 ; stem.gn.gamma
dq 3,64,64,1,1,1,0,0 ; stem.gn.beta
dq 1,36864,64,3,3,64,0,0 ; s0.b0.c1.w
dq 2,64,64,1,1,1,0,0 ; s0.b0.g1.gamma
dq 3,64,64,1,1,1,0,0 ; s0.b0.g1.beta
dq 1,36864,64,3,3,64,0,0 ; s0.b0.c2.w
dq 2,64,64,1,1,1,0,0 ; s0.b0.g2.gamma
dq 3,64,64,1,1,1,0,0 ; s0.b0.g2.beta
dq 1,36864,64,3,3,64,0,0 ; s0.b1.c1.w
dq 2,64,64,1,1,1,0,0 ; s0.b1.g1.gamma
dq 3,64,64,1,1,1,0,0 ; s0.b1.g1.beta
dq 1,36864,64,3,3,64,0,0 ; s0.b1.c2.w
dq 2,64,64,1,1,1,0,0 ; s0.b1.g2.gamma
dq 3,64,64,1,1,1,0,0 ; s0.b1.g2.beta
dq 1,73728,128,3,3,64,0,0 ; s1.b0.c1.w
dq 2,128,128,1,1,1,0,0 ; s1.b0.g1.gamma
dq 3,128,128,1,1,1,0,0 ; s1.b0.g1.beta
dq 1,147456,128,3,3,128,0,0 ; s1.b0.c2.w
dq 2,128,128,1,1,1,0,0 ; s1.b0.g2.gamma
dq 3,128,128,1,1,1,0,0 ; s1.b0.g2.beta
dq 1,8192,128,1,1,64,0,0 ; s1.b0.skip.w
dq 2,128,128,1,1,1,0,0 ; s1.b0.skipg.gamma
dq 3,128,128,1,1,1,0,0 ; s1.b0.skipg.beta
dq 1,147456,128,3,3,128,0,0 ; s1.b1.c1.w
dq 2,128,128,1,1,1,0,0 ; s1.b1.g1.gamma
dq 3,128,128,1,1,1,0,0 ; s1.b1.g1.beta
dq 1,147456,128,3,3,128,0,0 ; s1.b1.c2.w
dq 2,128,128,1,1,1,0,0 ; s1.b1.g2.gamma
dq 3,128,128,1,1,1,0,0 ; s1.b1.g2.beta
dq 1,294912,256,3,3,128,0,0 ; s2.b0.c1.w
dq 2,256,256,1,1,1,0,0 ; s2.b0.g1.gamma
dq 3,256,256,1,1,1,0,0 ; s2.b0.g1.beta
dq 1,589824,256,3,3,256,0,0 ; s2.b0.c2.w
dq 2,256,256,1,1,1,0,0 ; s2.b0.g2.gamma
dq 3,256,256,1,1,1,0,0 ; s2.b0.g2.beta
dq 1,32768,256,1,1,128,0,0 ; s2.b0.skip.w
dq 2,256,256,1,1,1,0,0 ; s2.b0.skipg.gamma
dq 3,256,256,1,1,1,0,0 ; s2.b0.skipg.beta
dq 1,589824,256,3,3,256,0,0 ; s2.b1.c1.w
dq 2,256,256,1,1,1,0,0 ; s2.b1.g1.gamma
dq 3,256,256,1,1,1,0,0 ; s2.b1.g1.beta
dq 1,589824,256,3,3,256,0,0 ; s2.b1.c2.w
dq 2,256,256,1,1,1,0,0 ; s2.b1.g2.gamma
dq 3,256,256,1,1,1,0,0 ; s2.b1.g2.beta
dq 4,2560,10,256,1,1,0,0 ; head.w
dq 5,10,10,1,1,1,0,0 ; head.b
size cfar_ops,CFAR_OP_COUNT*CFAR_OP_SIZE
PD_W equ 0
PD_MASTER equ 8
PD_GRAD equ 16
PD_M equ 24
PD_V equ 32
PD_EMA equ 40
PD_ELEMS equ 48
PD_KIND equ 56
PD_SIZE equ 64
section .bss align=64
cfar_gpu_params resb CFAR_GPU_MAX*CFAR_PARAM_TENSORS*PD_SIZE
cfar_gpu_tensors resb CFAR_GPU_MAX*CFAR_TENSOR_COUNT*T_SIZE
cfar_gpu_grads resb CFAR_GPU_MAX*CFAR_TENSOR_COUNT*T_SIZE
cfar_gpu_aux resb CFAR_GPU_MAX*CFAR_OP_COUNT*T_SIZE
cfar_stage_host resq CFAR_GPU_MAX*2
cfar_stage_gpu resq CFAR_GPU_MAX*2
cfar_labels_host resq CFAR_GPU_MAX*2
cfar_labels_gpu resq CFAR_GPU_MAX*2
cfar_loader_ready resd CFAR_GPU_MAX*2
cfar_loader_done resd CFAR_GPU_MAX*2
cfar_checkpoint_header resb CFAR_CKPT_HEADER
cfar_eval_confusion resq 100
cfar_eval_correct resq 1
cfar_eval_total resq 1
cfar_eval_nll_bits resq 1
section .text
cfar_param_desc:
; edi gpu,esi param; rax descriptor capability
mov eax,edi
imul rax,CFAR_PARAM_TENSORS
add rax,rsi
shl rax,6
lea rax,[rel cfar_gpu_params+rax]
mov rcx,PD_SIZE
bnd rax,rcx
ret
cfar_tensor_desc:
; edi gpu,esi tensor; rax
mov eax,edi
imul rax,CFAR_TENSOR_COUNT
add rax,rsi
shl rax,6
lea rax,[rel cfar_gpu_tensors+rax]
mov rcx,T_SIZE
bnd rax,rcx
ret
cfar_grad_desc:
mov eax,edi
imul rax,CFAR_TENSOR_COUNT
add rax,rsi
shl rax,6
lea rax,[rel cfar_gpu_grads+rax]
mov rcx,T_SIZE
bnd rax,rcx
ret
cfar_aux_desc:
mov eax,edi
imul rax,CFAR_OP_COUNT
add rax,rsi
shl rax,6
lea rax,[rel cfar_gpu_aux+rax]
mov rcx,T_SIZE
bnd rax,rcx
ret
cfar_open_gpus:
call gpu.count
cmp rax,CFAR_GPU_MAX
jb cfar_open_gpus_bad
xor ebx,ebx
cfar_open_gpus_loop:
cmp ebx,CFAR_GPU_MAX
je cfar_open_gpus_done
mov edi,ebx
mov esi,89
call gpu.open
test rax,rax
jz cfar_open_gpus_bad
mov qword [rel cfar_gpu_handles+rbx*8],rax
mov rdi,rax
call gpu.event_new
test rax,rax
jz cfar_open_gpus_bad
mov qword [rel cfar_gpu_events+rbx*8],rax
inc ebx
jmp cfar_open_gpus_loop
cfar_open_gpus_done:
ret
cfar_open_gpus_bad:
lea rdi,[rel cfar_err_gpu]
mov rsi,cfar_err_gpu_n
mov edx,69
jmp fatal
cfar_close_gpus:
xor ebx,ebx
cfar_close_gpus_loop:
cmp ebx,CFAR_GPU_MAX
je cfar_close_gpus_done
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
test rdi,rdi
jz cfar_close_gpus_next
mov rsi,qword [rel cfar_gpu_events+rbx*8]
call gpu.event_free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call gpu.close
mov qword [rel cfar_gpu_handles+rbx*8],0
cfar_close_gpus_next:
inc ebx
jmp cfar_close_gpus_loop
cfar_close_gpus_done:
ret
cfar_alloc_params:
; allocate six exact GPU tensors per parameter on each device
xor ebx,ebx
cfar_alloc_params_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_alloc_params_done
xor r12d,r12d
cfar_alloc_params_p:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_alloc_params_gpu_next
mov eax,r12d
shl rax,6
lea r13,[rel cfar_param_shapes+rax]
mov r14,qword [r13+8]
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r15,rax
mov qword [r15+PD_ELEMS],r14
mov rax,qword [r13]
mov qword [r15+PD_KIND],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*4]
call gpu.alloc
mov qword [r15+PD_MASTER],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*2]
call gpu.alloc
mov qword [r15+PD_W],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*4]
call gpu.alloc
mov qword [r15+PD_GRAD],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*4]
call gpu.alloc
mov qword [r15+PD_M],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*4]
call gpu.alloc
mov qword [r15+PD_V],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[r14*4]
call gpu.alloc
mov qword [r15+PD_EMA],rax
inc r12d
jmp cfar_alloc_params_p
cfar_alloc_params_gpu_next:
inc ebx
jmp cfar_alloc_params_gpu
cfar_alloc_params_done:
ret
cfar_alloc_graph:
; exact forward and gradient allocations for every graph tensor, local batch 128
xor ebx,ebx
cfar_alloc_graph_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_alloc_graph_done
xor r12d,r12d
cfar_alloc_graph_tensor:
cmp r12d,CFAR_TENSOR_COUNT
je cfar_alloc_graph_aux
mov eax,r12d
shl rax,5
lea r13,[rel cfar_tensor_shapes+rax]
mov r14,qword [r13+0]
imul r14,qword [r13+8]
imul r14,qword [r13+16]
imul r14,CFAR_LOCAL_BATCH
mov r15,2
cmp r12d,CFAR_TENSOR_COUNT-2
jb cfar_alloc_graph_type
mov r15,4
cfar_alloc_graph_type:
mov edi,ebx
mov esi,r12d
call cfar_tensor_desc
mov r10,rax
mov rsi,r10
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rdx,r14
imul rdx,r15
mov r8,DT_BF16
cmp r15,4
jne cfar_alloc_graph_fwd
mov r8,DT_F32
cfar_alloc_graph_fwd:
call tensor_gpu_new
mov rax,qword [r13+0]
mov qword [r10+T_D0],rax
mov rax,qword [r13+8]
mov qword [r10+T_D1],rax
mov rax,qword [r13+16]
mov qword [r10+T_D2],rax
mov qword [r10+T_D3],CFAR_LOCAL_BATCH
mov edi,ebx
mov esi,r12d
call cfar_grad_desc
mov r10,rax
mov rsi,r10
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rdx,r14
shl rdx,2
mov r8,DT_F32
call tensor_gpu_new
mov rax,qword [r13+0]
mov qword [r10+T_D0],rax
mov rax,qword [r13+8]
mov qword [r10+T_D1],rax
mov rax,qword [r13+16]
mov qword [r10+T_D2],rax
mov qword [r10+T_D3],CFAR_LOCAL_BATCH
inc r12d
jmp cfar_alloc_graph_tensor
cfar_alloc_graph_aux:
xor r12d,r12d
cfar_alloc_graph_aux_loop:
cmp r12d,CFAR_OP_COUNT
je cfar_alloc_graph_stage
mov edi,ebx
mov esi,r12d
call cfar_aux_desc
mov rsi,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rdx,CFAR_LOCAL_BATCH*32*2*4
mov r8,DT_F32
call tensor_gpu_new
inc r12d
jmp cfar_alloc_graph_aux_loop
cfar_alloc_graph_stage:
xor r15d,r15d
cfar_alloc_graph_stage_parity:
cmp r15d,2
je cfar_alloc_graph_stage_done
mov eax,ebx
shl eax,1
add eax,r15d
mov r14d,eax
mov rdi,CFAR_LOCAL_BATCH*CFAR_IMAGE_BYTES
call alloc
mov qword [rel cfar_stage_host+r14*8],rax
mov rdi,CFAR_LOCAL_BATCH
call alloc
mov qword [rel cfar_labels_host+r14*8],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,CFAR_LOCAL_BATCH*CFAR_IMAGE_BYTES
call gpu.alloc
mov qword [rel cfar_stage_gpu+r14*8],rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,CFAR_LOCAL_BATCH
call gpu.alloc
mov qword [rel cfar_labels_gpu+r14*8],rax
inc r15d
jmp cfar_alloc_graph_stage_parity
cfar_alloc_graph_stage_done:
inc ebx
jmp cfar_alloc_graph_gpu
cfar_alloc_graph_done:
ret
cfar_init_params:
; deterministic CLT-normal initialization on host, copied to every GPU, zeros optimizer
mov rdi,0x4f1bbcdc12345678
mov qword [rel cfar_rng_counter],rdi
xor r12d,r12d
cfar_init_params_loop:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_init_params_done
mov eax,r12d
shl rax,6
lea r13,[rel cfar_param_shapes+rax]
mov r14,qword [r13+8]
lea rdi,[r14*4]
call alloc
mov r15,rax
mov rdi,r15
mov rsi,r14
mov rdx,qword [r13]
mov rcx,r13
call cfar_fill_initial
xor ebx,ebx
cfar_init_params_copy_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_init_params_free
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r10,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r10+PD_MASTER]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r10+PD_EMA]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,r10
call cfar_cast_param_bf16
inc ebx
jmp cfar_init_params_copy_gpu
cfar_init_params_free:
mov rdi,r15
call free
inc r12d
jmp cfar_init_params_loop
cfar_init_params_done:
mov qword [rel cfar_step],0
mov qword [rel cfar_epoch],0
mov qword [rel cfar_loss_scale],CFAR_LOSS_SCALE_INIT
ret
cfar_fill_initial:
; rdi f32 dst,rsi elements,rdx kind,rcx shape descriptor
; GN gamma=1, beta/bias=0, weights CLT normal scaled by fan-in.
push rbx
push r12
push r13
push r14
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov r14,rcx
cmp r13,2
je cfar_fill_ones
cmp r13,3
je cfar_fill_zero
cmp r13,5
je cfar_fill_zero
; scale=sqrt(2/fan_in), fan_in product of shape dimensions 1..3 for conv, dim1 for linear.
mov rax,qword [r14+24]
imul rax,qword [r14+32]
imul rax,qword [r14+40]
cmp r13,4
jne cfar_fill_fan_ok
mov rax,qword [r14+24]
cfar_fill_fan_ok:
cvtsi2ss xmm0,rax
movss xmm1,dword [rel f_two]
divss xmm1,xmm0
sqrtss xmm7,xmm1
xor r8d,r8d
cfar_fill_random_loop:
cmp r8,r12
je cfar_fill_done
xorps xmm0,xmm0
mov ecx,12
cfar_fill_random_sum:
call cfar_rng_u32
cvtsi2ss xmm1,eax
mulss xmm1,dword [rel f_inv_u32]
addss xmm0,xmm1
dec ecx
jnz cfar_fill_random_sum
subss xmm0,dword [rel f_six]
mulss xmm0,xmm7
movss dword [rbx+r8*4],xmm0
inc r8
jmp cfar_fill_random_loop
cfar_fill_ones:
movss xmm0,dword [rel f_one]
xor r8d,r8d
cfar_fill_ones_loop:
cmp r8,r12
je cfar_fill_done
movss dword [rbx+r8*4],xmm0
inc r8
jmp cfar_fill_ones_loop
cfar_fill_zero:
mov rdi,rbx
lea rsi,[r12*4]
call memzero
cfar_fill_done:
pop r14
pop r13
pop r12
pop rbx
ret
cfar_rng_u32:
; splitmix64 deterministic; eax output
mov rax,qword [rel cfar_rng_counter]
add rax,0x9e3779b97f4a7c15
mov qword [rel cfar_rng_counter],rax
mov rdx,rax
shr rdx,30
xor rax,rdx
mov rdx,0xbf58476d1ce4e5b9
imul rax,rdx
mov rdx,rax
shr rdx,27
xor rax,rdx
mov rdx,0x94d049bb133111eb
imul rax,rdx
mov rdx,rax
shr rdx,31
xor rax,rdx
shr rax,32
ret
section .rodata
f_six dd 6.0
section .text
launch_1d:
; rdi device,rsi kernel,rdx args,rcx argc,r8 total,r9 block
push rbx
push r12
push r13
push r14
lea rsp,[rsp-64]
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov r14,rcx
mov rax,r8
lea rax,[rax+r9-1]
xor edx,edx
div r9
mov edi,eax
mov esi,1
mov edx,1
call gpu_pack_xyz
mov qword [rsp+LD_GRID],rax
mov edi,r9d
mov esi,1
mov edx,1
call gpu_pack_xyz
mov qword [rsp+LD_BLOCK],rax
mov qword [rsp+LD_SMEM],0
mov qword [rsp+LD_ARGS],r13
mov qword [rsp+LD_ARGC],r14
mov rdi,rbx
mov rsi,r12
mov rdx,rsp
call gpu_launch_desc
lea rsp,[rsp+64]
pop r14
pop r13
pop r12
pop rbx
ret
cfar_cast_param_bf16:
; rdi device,rsi parameter descriptor
push rbx
push r12
lea rsp,[rsp-160]
mov rbx,rdi
mov r12,rsi
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r12+PD_MASTER]
xor edx,edx
mov rcx,qword [r12+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r12+PD_W]
xor edx,edx
mov rcx,qword [r12+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r12+PD_ELEMS]
call gpu_arg_u64
mov rdi,rbx
lea rsi,[rel cfar_k_f32_to_bf16]
mov rdx,rsp
mov rcx,3
mov r8,qword [r12+PD_ELEMS]
mov r9,256
call launch_1d
lea rsp,[rsp+160]
pop r12
pop rbx
ret
cfar_zero_grads:
xor ebx,ebx
cfar_zero_grads_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_zero_grads_done
xor r12d,r12d
cfar_zero_grads_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_zero_grads_gpu_next
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r13,rax
lea rsp,[rsp-96]
lea rdi,[rsp]
mov rsi,qword [r13+PD_GRAD]
xor edx,edx
mov rcx,qword [r13+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+GA_SIZE]
mov rsi,qword [r13+PD_ELEMS]
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_zero_f32]
mov rdx,rsp
mov rcx,2
mov r8,qword [r13+PD_ELEMS]
mov r9,256
call launch_1d
lea rsp,[rsp+96]
inc r12d
jmp cfar_zero_grads_param
cfar_zero_grads_gpu_next:
inc ebx
jmp cfar_zero_grads_gpu
cfar_zero_grads_done:
ret
cfar_forward:
; edi gpu
push rbx
push r12
mov ebx,edi
xor r12d,r12d
cfar_forward_loop:
cmp r12d,CFAR_OP_COUNT-1
jae cfar_forward_done
mov eax,r12d
shl rax,6
lea rsi,[rel cfar_ops+rax]
mov edi,ebx
call cfar_forward_op
test eax,eax
jnz cfar_forward_bad
inc r12d
jmp cfar_forward_loop
cfar_forward_done:
xor eax,eax
pop r12
pop rbx
ret
cfar_forward_bad:
mov eax,-1
pop r12
pop rbx
ret
cfar_forward_op:
; edi gpu,rsi op descriptor
mov eax,dword [rsi+OD_KIND]
cmp eax,OP_CONV
je cfar_launch_conv_fwd
cmp eax,OP_GN
je cfar_launch_gn_fwd
cmp eax,OP_SILU
je cfar_launch_silu_fwd
cmp eax,OP_ADD
je cfar_launch_add_fwd
cmp eax,OP_GAP
je cfar_launch_gap_fwd
cmp eax,OP_LINEAR
je cfar_launch_linear_fwd
mov eax,-1
ret
cfar_launch_conv_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-512]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r14,rax
mov esi,dword [r12+OD_PARAM]
mov edi,ebx
call cfar_param_desc
mov r15,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r15+PD_W]
xor edx,edx
mov rcx,qword [r15+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
mov eax,dword [r12+OD_IN0]
shl rax,5
lea r10,[rel cfar_tensor_shapes+rax]
mov eax,dword [r12+OD_OUT]
shl rax,5
lea r11,[rel cfar_tensor_shapes+rax]
lea rdi,[rsp+3*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+4*GA_SIZE]
mov rsi,qword [r10+0]
call gpu_arg_u64
lea rdi,[rsp+5*GA_SIZE]
mov rsi,qword [r10+8]
call gpu_arg_u64
lea rdi,[rsp+6*GA_SIZE]
mov rsi,qword [r10+16]
call gpu_arg_u64
lea rdi,[rsp+7*GA_SIZE]
mov rsi,qword [r11+16]
call gpu_arg_u64
mov rax,qword [r12+OD_X0]
movzx r9d,al
shr rax,8
movzx r10d,al
shr rax,8
movzx r11d,al
lea rdi,[rsp+8*GA_SIZE]
mov rsi,r9
call gpu_arg_u64
lea rdi,[rsp+9*GA_SIZE]
mov rsi,r10
call gpu_arg_u64
lea rdi,[rsp+10*GA_SIZE]
mov rsi,r11
call gpu_arg_u64
mov r8,CFAR_LOCAL_BATCH
imul r8,qword [r14+T_D0]
imul r8,qword [r14+T_D1]
imul r8,qword [r14+T_D2]
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_conv_fwd]
mov rdx,rsp
mov rcx,11
mov r9,256
call launch_1d
lea rsp,[rsp+512]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_launch_gn_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-640]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r14,rax
mov esi,dword [r12+OD_PARAM]
mov edi,ebx
call cfar_param_desc
mov r15,rax
mov eax,dword [r12+OD_AUX]
mov esi,eax
mov edi,ebx
call cfar_param_desc
mov r10,rax
mov rax,r12
lea rcx,[rel cfar_ops]
sub rax,rcx
shr rax,6
mov qword [rsp+608],rax
mov esi,eax
mov edi,ebx
call cfar_aux_desc
mov r11,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r15+PD_W]
xor edx,edx
mov rcx,qword [r15+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r10+PD_W]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+4*GA_SIZE]
mov rsi,qword [r11+T_CAP]
xor edx,edx
mov rcx,qword [r11+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
mov eax,dword [r12+OD_IN0]
shl rax,5
lea r9,[rel cfar_tensor_shapes+rax]
lea rdi,[rsp+5*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+6*GA_SIZE]
mov rsi,qword [r9+0]
call gpu_arg_u64
lea rdi,[rsp+7*GA_SIZE]
mov rsi,qword [r9+8]
call gpu_arg_u64
lea rdi,[rsp+8*GA_SIZE]
mov rsi,qword [r9+16]
call gpu_arg_u64
lea rdi,[rsp+9*GA_SIZE]
mov rsi,qword [r12+OD_X0]
call gpu_arg_u64
mov r8,CFAR_LOCAL_BATCH
imul r8,qword [r12+OD_X0]
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_gn_fwd]
mov rdx,rsp
mov rcx,10
mov r9,256
call launch_1d
lea rsp,[rsp+640]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_launch_silu_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
lea rsp,[rsp-160]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r10,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
mov r8,qword [r10+T_BYTES]
shr r8,1
lea rdi,[rsp+2*GA_SIZE]
mov rsi,r8
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_silu_fwd]
mov rdx,rsp
mov rcx,3
mov r9,256
call launch_1d
lea rsp,[rsp+160]
pop r13
pop r12
pop rbx
ret
cfar_launch_add_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
push r14
lea rsp,[rsp-192]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_IN1]
mov edi,ebx
call cfar_tensor_desc
mov r14,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r10,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
mov r8,qword [r10+T_BYTES]
shr r8,1
lea rdi,[rsp+3*GA_SIZE]
mov rsi,r8
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_add_fwd]
mov rdx,rsp
mov rcx,4
mov r9,256
call launch_1d
lea rsp,[rsp+192]
pop r14
pop r13
pop r12
pop rbx
ret
cfar_launch_gap_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
lea rsp,[rsp-256]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r10,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
mov eax,dword [r12+OD_IN0]
shl rax,5
lea r11,[rel cfar_tensor_shapes+rax]
lea rdi,[rsp+2*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [r11+0]
call gpu_arg_u64
lea rdi,[rsp+4*GA_SIZE]
mov rsi,qword [r11+8]
call gpu_arg_u64
lea rdi,[rsp+5*GA_SIZE]
mov rsi,qword [r11+16]
call gpu_arg_u64
mov r8,CFAR_LOCAL_BATCH
imul r8,qword [r11+16]
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_gap_fwd]
mov rdx,rsp
mov rcx,6
mov r9,256
call launch_1d
lea rsp,[rsp+256]
pop r13
pop r12
pop rbx
ret
cfar_launch_linear_fwd:
; edi gpu,rsi op
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-320]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r14,rax
mov esi,dword [r12+OD_PARAM]
mov edi,ebx
call cfar_param_desc
mov r15,rax
mov esi,dword [r12+OD_AUX]
mov edi,ebx
call cfar_param_desc
mov r10,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r15+PD_W]
xor edx,edx
mov rcx,qword [r15+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r10+PD_W]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+4*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+5*GA_SIZE]
mov rsi,qword [r12+OD_X0]
call gpu_arg_u64
lea rdi,[rsp+6*GA_SIZE]
mov rsi,qword [r12+OD_X1]
call gpu_arg_u64
mov r8,CFAR_LOCAL_BATCH
imul r8,qword [r12+OD_X1]
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_linear_fwd]
mov rdx,rsp
mov rcx,7
mov r9,256
call launch_1d
lea rsp,[rsp+320]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_backward:
; edi gpu
push rbx
push r12
mov ebx,edi
mov r12d,CFAR_OP_COUNT-2
cfar_backward_loop:
cmp r12d,-1
je cfar_backward_done
mov eax,r12d
shl rax,6
lea rsi,[rel cfar_ops+rax]
mov edi,ebx
call cfar_backward_op
test eax,eax
jnz cfar_backward_bad
dec r12d
jmp cfar_backward_loop
cfar_backward_done:
xor eax,eax
pop r12
pop rbx
ret
cfar_backward_bad:
mov eax,-1
pop r12
pop rbx
ret
cfar_backward_op:
; edi gpu,rsi op; common 16-argument launch
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-640]
mov ebx,edi
mov r12,rsi
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_tensor_desc
mov r13,rax
mov esi,dword [r12+OD_IN0]
mov edi,ebx
call cfar_grad_desc
mov r14,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_tensor_desc
mov r15,rax
mov esi,dword [r12+OD_OUT]
mov edi,ebx
call cfar_grad_desc
mov qword [rsp+560],rax
mov esi,dword [r12+OD_IN1]
mov edi,ebx
call cfar_tensor_desc
mov qword [rsp+568],rax
mov esi,dword [r12+OD_IN1]
mov edi,ebx
call cfar_grad_desc
mov qword [rsp+576],rax
mov esi,dword [r12+OD_PARAM]
mov edi,ebx
call cfar_param_desc
mov qword [rsp+584],rax
mov esi,dword [r12+OD_AUX]
mov edi,ebx
call cfar_param_desc
mov qword [rsp+592],rax
mov rax,r12
lea rcx,[rel cfar_ops]
sub rax,rcx
shr rax,6
mov esi,eax
mov edi,ebx
call cfar_aux_desc
mov qword [rsp+600],rax
; x,y,dy,dx,x2,dx2,w,dw,p2,dp2,aux
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r15+T_CAP]
xor edx,edx
mov rcx,qword [r15+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
mov r10,qword [rsp+560]
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_RW
call gpu_arg_span
mov r10,qword [rsp+568]
lea rdi,[rsp+4*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
mov r10,qword [rsp+576]
lea rdi,[rsp+5*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_RW
call gpu_arg_span
mov r10,qword [rsp+584]
lea rdi,[rsp+6*GA_SIZE]
mov rsi,qword [r10+PD_W]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+7*GA_SIZE]
mov rsi,qword [r10+PD_GRAD]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
mov r10,qword [rsp+592]
lea rdi,[rsp+8*GA_SIZE]
mov rsi,qword [r10+PD_W]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+9*GA_SIZE]
mov rsi,qword [r10+PD_GRAD]
xor edx,edx
mov rcx,qword [r10+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
mov r10,qword [rsp+600]
lea rdi,[rsp+10*GA_SIZE]
mov rsi,qword [r10+T_CAP]
xor edx,edx
mov rcx,qword [r10+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+11*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+12*GA_SIZE]
mov rsi,qword [r12+OD_X0]
call gpu_arg_u64
lea rdi,[rsp+13*GA_SIZE]
mov rsi,qword [r12+OD_X1]
call gpu_arg_u64
lea rdi,[rsp+14*GA_SIZE]
mov rsi,qword [rsp+608]
call gpu_arg_u64
lea rdi,[rsp+15*GA_SIZE]
mov rsi,qword [rel cfar_loss_scale]
call gpu_arg_u64
mov rax,qword [r13+T_BYTES]
shr rax,1
mov r8,rax
mov rax,qword [r15+T_BYTES]
shr rax,1
add r8,rax
mov r10,qword [rsp+584]
add r8,qword [r10+PD_ELEMS]
mov rax,qword [r12+OD_KIND]
cmp rax,OP_CONV
je cfar_backward_kernel_conv
cmp rax,OP_GN
je cfar_backward_kernel_gn
cmp rax,OP_SILU
je cfar_backward_kernel_silu
cmp rax,OP_ADD
je cfar_backward_kernel_add
cmp rax,OP_GAP
je cfar_backward_kernel_gap
cmp rax,OP_LINEAR
je cfar_backward_kernel_linear
mov eax,-1
jmp cfar_backward_ret
cfar_backward_kernel_conv:
lea rsi,[rel cfar_k_conv_bwd]
mov rax,qword [r13+T_BYTES]
shr rax,1
mov r8,rax
mov r10,qword [rsp+584]
add r8,qword [r10+PD_ELEMS]
jmp cfar_backward_launch
cfar_backward_kernel_gn:
lea rsi,[rel cfar_k_gn_bwd]
mov r8,CFAR_LOCAL_BATCH
imul r8,qword [r12+OD_X0]
jmp cfar_backward_launch
cfar_backward_kernel_silu:
lea rsi,[rel cfar_k_silu_bwd]
mov r8,qword [r15+T_BYTES]
shr r8,1
jmp cfar_backward_launch
cfar_backward_kernel_add:
lea rsi,[rel cfar_k_add_bwd]
mov r8,qword [r15+T_BYTES]
shr r8,1
jmp cfar_backward_launch
cfar_backward_kernel_gap:
lea rsi,[rel cfar_k_gap_bwd]
mov r8,qword [r13+T_BYTES]
shr r8,1
jmp cfar_backward_launch
cfar_backward_kernel_linear:
lea rsi,[rel cfar_k_linear_bwd]
mov rax,qword [r13+T_BYTES]
shr rax,1
mov r8,rax
mov r10,qword [rsp+584]
add r8,qword [r10+PD_ELEMS]
mov r10,qword [rsp+592]
add r8,qword [r10+PD_ELEMS]
cfar_backward_launch:
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rdx,rsp
mov rcx,16
mov r9,256
call launch_1d
cfar_backward_ret:
lea rsp,[rsp+640]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_record_ptr:
; rdi train index 0..49999; rax record authority exact 3073
xor edx,edx
mov rax,rdi
mov rcx,10000
div rcx
mov r8,qword [rel cfar_batch_maps+rax*8]
imul rdx,CFAR_RECORD_BYTES
lea rax,[r8+rdx]
mov rcx,CFAR_RECORD_BYTES
bnd rax,rcx
ret
cfar_feistel:
; edi value <65536, esi epoch; eax permutation over 16 bits
mov eax,edi
mov r8d,esi
mov ecx,4
cfar_feistel_round:
mov edx,eax
and edx,0xff
shr eax,8
mov r9d,edx
imul r9d,r9d,0x9d
add r9d,r8d
rol r9d,cl
xor r9d,eax
and r9d,0xff
shl edx,8
or edx,r9d
mov eax,edx
add r8d,0x3d
loop cfar_feistel_round
ret
cfar_permute_index:
; edi position 0..49999,esi epoch; eax unique 0..49999
mov r10d,edi
cfar_permute_cycle:
mov edi,r10d
call cfar_feistel
cmp eax,CFAR_TRAIN_IMAGES
jb cfar_permute_done
mov r10d,eax
jmp cfar_permute_cycle
cfar_permute_done:
ret
cfar_fill_stage:
; edi gpu,esi step,edx parity,ecx eval(0 train,1 test)
push rbx
push r12
push r13
push r14
push r15
mov ebx,edi
mov r12d,esi
mov r13d,edx
mov r14d,ecx
mov eax,ebx
shl rax,1
add eax,r13d
mov r15,qword [rel cfar_stage_host+rax*8]
mov r11,qword [rel cfar_labels_host+rax*8]
xor r10d,r10d
cfar_fill_stage_loop:
cmp r10d,CFAR_LOCAL_BATCH
je cfar_fill_stage_done
cmp r10d,CFAR_VALID_PER_GPU
jae cfar_fill_stage_pad
test r14d,r14d
jnz cfar_fill_stage_test
mov eax,r12d
imul eax,CFAR_GLOBAL_BATCH
mov ecx,ebx
imul ecx,CFAR_VALID_PER_GPU
add eax,ecx
add eax,r10d
mov edi,eax
mov esi,dword [rel cfar_epoch]
call cfar_permute_index
mov edi,eax
call cfar_record_ptr
jmp cfar_fill_stage_copy
cfar_fill_stage_test:
mov eax,r12d
imul eax,CFAR_GLOBAL_BATCH
mov ecx,ebx
imul ecx,CFAR_VALID_PER_GPU
add eax,ecx
add eax,r10d
cmp eax,CFAR_TEST_IMAGES
jae cfar_fill_stage_pad
mov r8,qword [rel cfar_batch_maps+5*8]
imul rax,CFAR_RECORD_BYTES
lea rax,[r8+rax]
mov rcx,CFAR_RECORD_BYTES
bnd rax,rcx
cfar_fill_stage_copy:
mov dl,byte [rax]
mov byte [r11+r10],dl
lea rsi,[rax+1]
mov rax,r10
imul rax,CFAR_IMAGE_BYTES
lea rdi,[r15+rax]
mov rdx,CFAR_IMAGE_BYTES
call memcpy_data
inc r10d
jmp cfar_fill_stage_loop
cfar_fill_stage_pad:
mov byte [r11+r10],255
mov rax,r10
imul rax,CFAR_IMAGE_BYTES
lea rdi,[r15+rax]
mov rsi,CFAR_IMAGE_BYTES
call memzero
inc r10d
jmp cfar_fill_stage_loop
cfar_fill_stage_done:
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_copy_stage_gpu:
; edi gpu,esi parity
push rbx
mov ebx,edi
mov eax,ebx
shl rax,1
add eax,esi
mov r10,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_stage_gpu+r10*8]
mov rdx,qword [rel cfar_stage_host+r10*8]
mov rcx,CFAR_LOCAL_BATCH*CFAR_IMAGE_BYTES
call gpu.copy_h2d
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_labels_gpu+r10*8]
mov rdx,qword [rel cfar_labels_host+r10*8]
mov rcx,CFAR_LOCAL_BATCH
call gpu.copy_h2d
pop rbx
ret
cfar_preprocess:
; edi gpu,esi parity,edx training
push rbx
push r12
push r13
lea rsp,[rsp-512]
mov ebx,edi
mov r12d,esi
mov r13d,edx
mov eax,ebx
shl rax,1
add eax,r12d
mov r10,rax
mov edi,ebx
xor esi,esi
call cfar_tensor_desc
mov r11,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [rel cfar_stage_gpu+r10*8]
xor edx,edx
mov rcx,CFAR_LOCAL_BATCH*CFAR_IMAGE_BYTES
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r11+T_CAP]
xor edx,edx
mov rcx,qword [r11+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+3*GA_SIZE]
mov esi,dword [rel cfar_epoch]
call gpu_arg_u64
lea rdi,[rsp+4*GA_SIZE]
mov esi,dword [rel cfar_step]
call gpu_arg_u64
lea rdi,[rsp+5*GA_SIZE]
mov esi,r13d
call gpu_arg_u64
lea rdi,[rsp+6*GA_SIZE]
mov esi,dword [rel cfar_mean+0]
call gpu_arg_f32
lea rdi,[rsp+7*GA_SIZE]
mov esi,dword [rel cfar_mean+4]
call gpu_arg_f32
lea rdi,[rsp+8*GA_SIZE]
mov esi,dword [rel cfar_mean+8]
call gpu_arg_f32
lea rdi,[rsp+9*GA_SIZE]
mov esi,dword [rel cfar_invstd+0]
call gpu_arg_f32
lea rdi,[rsp+10*GA_SIZE]
mov esi,dword [rel cfar_invstd+4]
call gpu_arg_f32
lea rdi,[rsp+11*GA_SIZE]
mov esi,dword [rel cfar_invstd+8]
call gpu_arg_f32
mov r8,CFAR_LOCAL_BATCH*CFAR_IMAGE_BYTES
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_preprocess]
mov rdx,rsp
mov rcx,12
mov r9,256
call launch_1d
lea rsp,[rsp+512]
pop r13
pop r12
pop rbx
ret
cfar_zero_tensor_grads:
xor ebx,ebx
cfar_zero_tensor_grads_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_zero_tensor_grads_done
xor r12d,r12d
cfar_zero_tensor_grads_t:
cmp r12d,CFAR_TENSOR_COUNT
je cfar_zero_tensor_grads_next
mov edi,ebx
mov esi,r12d
call cfar_grad_desc
mov r13,rax
lea rsp,[rsp-96]
lea rdi,[rsp]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+GA_SIZE]
mov rsi,qword [r13+T_BYTES]
shr rsi,2
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_zero_f32]
mov rdx,rsp
mov rcx,2
mov r8,qword [r13+T_BYTES]
shr r8,2
mov r9,256
call launch_1d
lea rsp,[rsp+96]
inc r12d
jmp cfar_zero_tensor_grads_t
cfar_zero_tensor_grads_next:
inc ebx
jmp cfar_zero_tensor_grads_gpu
cfar_zero_tensor_grads_done:
ret
cfar_xent:
; edi gpu,esi parity,edx training
push rbx
push r12
lea rsp,[rsp-320]
mov ebx,edi
mov r12d,esi
mov dword [rsp+304],edx
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-2
call cfar_tensor_desc
mov r13,rax
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-2
call cfar_grad_desc
mov r14,rax
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-1
call cfar_tensor_desc
mov r15,rax
mov eax,ebx
shl rax,1
add eax,r12d
mov r10,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [rel cfar_labels_gpu+r10*8]
xor edx,edx
mov rcx,CFAR_LOCAL_BATCH
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r14+T_CAP]
xor edx,edx
mov rcx,qword [r14+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [r15+T_CAP]
xor edx,edx
mov rcx,qword [r15+T_BYTES]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+4*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
lea rdi,[rsp+5*GA_SIZE]
cmp dword [rsp+304],0
je cfar_xent_eval_smooth
mov esi,CFAR_LABEL_SMOOTH_BITS
jmp cfar_xent_set_smooth
cfar_xent_eval_smooth:
xor esi,esi
cfar_xent_set_smooth:
call gpu_arg_f32
lea rdi,[rsp+6*GA_SIZE]
cmp dword [rsp+304],0
je cfar_xent_eval_scale
mov rsi,qword [rel cfar_loss_scale]
jmp cfar_xent_set_scale
cfar_xent_eval_scale:
mov rsi,1
cfar_xent_set_scale:
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_xent]
mov rdx,rsp
mov rcx,7
mov r8,CFAR_LOCAL_BATCH
mov r9,64
call launch_1d
lea rsp,[rsp+320]
pop r12
pop rbx
ret
; ---- complete training, checkpoint, evaluation, and service control ----
CFAR_PARAM_ELEMS equ 2777674
CFAR_MAX_PARAM_ELEMS equ 589824
CFAR_PAYLOAD_BYTES equ CFAR_PARAM_ELEMS*16
CFAR_CK_VERSION equ 3
CFAR_CK_MAGIC equ 0
CFAR_CK_VERSION_O equ 16
CFAR_CK_HEADER_O equ 20
CFAR_CK_PAYLOAD_O equ 24
CFAR_CK_PARAM_O equ 32
CFAR_CK_STEP_O equ 40
CFAR_CK_EPOCH_O equ 48
CFAR_CK_LOSS_SCALE_O equ 56
CFAR_CK_RNG_O equ 64
CFAR_CK_MEAN_O equ 96
CFAR_CK_INVSTD_O equ 108
CFAR_CK_ARCH_O equ 128
CFAR_CK_DATA_O equ 160
CFAR_CK_PAYLOAD_SHA_O equ 192
CFAR_CK_HEADER_SHA_O equ 224
CFAR_CK_CREATED_NS_O equ 256
CFAR_CK_FLAGS_O equ 264
CFAR_WORK_ARG_SIZE equ 64
CFAR_WORKERS equ 96
CFAR_SERVE_BUF equ 65536
CFAR_SLOT_SIZE equ 69632
CFAR_SLOT_STATE equ 0
CFAR_SLOT_FD equ 8
CFAR_SLOT_GPU equ 16
CFAR_SLOT_REQ equ 64
CFAR_SLOT_META equ CFAR_SLOT_REQ+CFAR_SERVE_BUF
CFAR_SLOT_LOGITS equ CFAR_SLOT_META+64
CFAR_SLOT_JSON equ CFAR_SLOT_LOGITS+64
section .rodata align=64
cfar_dataset_sha db 0xc4,0xa3,0x8c,0x50,0xa1,0xbc,0x5f,0x3a,0x1c,0x55,0x37,0xf2,0x15,0x5a,0xb9,0xd6,0x8f,0x9f,0x25,0xeb,0x1e,0xd8,0xd9,0xdd,0xda,0x3d,0xb2,0x9a,0x59,0xbc,0xa1,0xdd
cfar_ck_magic16 db "WOWCFAR3-CHECK",0,0
cfar_json_prefix db "{\"class\":",0
cfar_json_name db ",\"name\":\"",0
cfar_json_probs db "\",\"probabilities\":[",0
cfar_json_tail db "]}",10,0
cfar_eval_prefix db "{\"test_images\":10000,\"correct\":",0
cfar_eval_mid db ",\"accuracy\":",0
cfar_eval_nll db ",\"nll\":",0
cfar_eval_end db "}",10,0
cfar_health_headers db "HTTP/1.1 200 OK",13,10,"Connection: close",13,10,"Content-Type: application/json",13,10,"Content-Length: 36",13,10,13,10
cfar_health_headers_n equ $-cfar_health_headers
cfar_beta1 dd 0.9
cfar_beta2 dd 0.999
cfar_adam_eps dd 1.0e-8
cfar_ema_decay dd 0.9999
cfar_ckpt_content_flags dq 0x000000000000000f
section .bss align=64
cfar_runtime_live resd 1
cfar_reduce_scratch resq 1
cfar_finite_flag_gpu resq CFAR_GPU_MAX
cfar_finite_flag_host resd CFAR_GPU_MAX
cfar_stats_total resq 6
cfar_ck_payload resq 1
cfar_stats_args resb CFAR_WORKERS*CFAR_WORK_ARG_SIZE
cfar_stats_threads resq CFAR_WORKERS
cfar_stats_sum resq CFAR_WORKERS*3
cfar_stats_sq resq CFAR_WORKERS*3
cfar_eval_gpu resq CFAR_GPU_MAX
cfar_eval_host resq CFAR_GPU_MAX
cfar_serve_fd resd 1
cfar_serve_threads resq CFAR_SERVE_WORKERS
cfar_serve_args resb CFAR_SERVE_WORKERS*16
cfar_gpu_locks resd CFAR_GPU_MAX
cfar_serve_storage resb CFAR_SERVE_WORKERS*CFAR_SLOT_SIZE
cfar_tmp_sha_ctx resb 112
cfar_tmp_sha resb 32
cfar_tmp_header_sha resb 32
cfar_float_text resb 64
cfar_eval_nll_f64 resq 1
section .text
cfar_runtime_open:
cmp dword [rel cfar_runtime_live],0
jne cfar_runtime_open_done
call cfar_batches_valid
test eax,eax
jz cfar_runtime_open_bad
call cfar_map_batches
call cfar_compute_stats
call cfar_open_gpus
call cfar_alloc_params
call cfar_alloc_graph
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,CFAR_MAX_PARAM_ELEMS*4
call gpu.alloc
mov qword [rel cfar_reduce_scratch],rax
xor ebx,ebx
cfar_runtime_flag_alloc:
cmp ebx,CFAR_GPU_MAX
je cfar_runtime_eval_alloc_start
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,4
call gpu.alloc
mov qword [rel cfar_finite_flag_gpu+rbx*8],rax
inc ebx
jmp cfar_runtime_flag_alloc
cfar_runtime_eval_alloc_start:
xor ebx,ebx
cfar_runtime_eval_alloc:
cmp ebx,CFAR_GPU_MAX
je cfar_runtime_open_mark
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,CFAR_LOCAL_BATCH*16
call gpu.alloc
mov qword [rel cfar_eval_gpu+rbx*8],rax
mov rdi,CFAR_LOCAL_BATCH*16
call alloc
mov qword [rel cfar_eval_host+rbx*8],rax
inc ebx
jmp cfar_runtime_eval_alloc
cfar_runtime_open_mark:
mov dword [rel cfar_runtime_live],1
cfar_runtime_open_done:
ret
cfar_runtime_open_bad:
lea rdi,[rel cfar_err_data]
mov rsi,cfar_err_data_n
mov edx,65
jmp fatal
cfar_compute_stats:
; exact byte sums and squared sums; 96 deterministic disjoint workers.
lea rdi,[rel cfar_stats_sum]
mov rsi,CFAR_WORKERS*3*8
call memzero
lea rdi,[rel cfar_stats_sq]
mov rsi,CFAR_WORKERS*3*8
call memzero
xor ebx,ebx
cfar_compute_stats_spawn:
cmp ebx,CFAR_WORKERS
je cfar_compute_stats_join
mov eax,CFAR_TRAIN_IMAGES
mul ebx
xor edx,edx
mov ecx,CFAR_WORKERS
div ecx
mov r12d,eax
mov eax,CFAR_TRAIN_IMAGES
lea ecx,[rbx+1]
mul ecx
xor edx,edx
mov ecx,CFAR_WORKERS
div ecx
mov r13d,eax
mov eax,ebx
shl rax,6
lea r14,[rel cfar_stats_args+rax]
mov dword [r14+0],r12d
mov dword [r14+4],r13d
lea rax,[rel cfar_stats_sum+rbx*24]
mov qword [r14+8],rax
lea rax,[rel cfar_stats_sq+rbx*24]
mov qword [r14+16],rax
lea rdi,[rel cfar_stats_worker]
mov rsi,r14
call os.thread_start
mov qword [rel cfar_stats_threads+rbx*8],rax
inc ebx
jmp cfar_compute_stats_spawn
cfar_compute_stats_join:
xor ebx,ebx
cfar_compute_stats_join_loop:
cmp ebx,CFAR_WORKERS
je cfar_compute_stats_reduce
mov rdi,qword [rel cfar_stats_threads+rbx*8]
call os.thread_join
inc ebx
jmp cfar_compute_stats_join_loop
cfar_compute_stats_reduce:
lea rdi,[rel cfar_stats_total]
mov rsi,48
call memzero
xor ebx,ebx
cfar_compute_stats_reduce_loop:
cmp ebx,CFAR_WORKERS
je cfar_compute_stats_finish
mov rax,qword [rel cfar_stats_sum+rbx*24+0]
add qword [rel cfar_stats_total+0],rax
mov rax,qword [rel cfar_stats_sum+rbx*24+8]
add qword [rel cfar_stats_total+8],rax
mov rax,qword [rel cfar_stats_sum+rbx*24+16]
add qword [rel cfar_stats_total+16],rax
mov rax,qword [rel cfar_stats_sq+rbx*24+0]
add qword [rel cfar_stats_total+24],rax
mov rax,qword [rel cfar_stats_sq+rbx*24+8]
add qword [rel cfar_stats_total+32],rax
mov rax,qword [rel cfar_stats_sq+rbx*24+16]
add qword [rel cfar_stats_total+40],rax
inc ebx
jmp cfar_compute_stats_reduce_loop
cfar_compute_stats_finish:
mov r14,CFAR_TRAIN_IMAGES*1024
cvtsi2ss xmm7,r14
mov rax,qword [rel cfar_stats_total+0]
mov rdx,qword [rel cfar_stats_total+24]
call cfar_stats_one
movss dword [rel cfar_mean+0],xmm0
movss dword [rel cfar_invstd+0],xmm1
mov rax,qword [rel cfar_stats_total+8]
mov rdx,qword [rel cfar_stats_total+32]
call cfar_stats_one
movss dword [rel cfar_mean+4],xmm0
movss dword [rel cfar_invstd+4],xmm1
mov rax,qword [rel cfar_stats_total+16]
mov rdx,qword [rel cfar_stats_total+40]
call cfar_stats_one
movss dword [rel cfar_mean+8],xmm0
movss dword [rel cfar_invstd+8],xmm1
ret
cfar_stats_one:
; rax sum,rdx sumsq,xmm7 count; xmm0 mean[0,1],xmm1 inverse std.
cvtsi2ss xmm0,rax
divss xmm0,xmm7
movss xmm2,xmm0
movss xmm3,dword [rel f_inv_255]
mulss xmm0,xmm3
cvtsi2ss xmm1,rdx
divss xmm1,xmm7
mulss xmm2,xmm2
subss xmm1,xmm2
maxss xmm1,dword [rel f_one_4096]
sqrtss xmm1,xmm1
movss xmm2,dword [rel f_255]
divss xmm2,xmm1
movss xmm1,xmm2
ret
cfar_stats_worker:
; rdi points to exact 64-byte argument.
push rbx
push r12
push r13
push r14
push r15
mov rbx,rdi
mov r12d,dword [rbx+0]
mov r13d,dword [rbx+4]
mov r14,qword [rbx+8]
mov r15,qword [rbx+16]
cfar_stats_worker_record:
cmp r12d,r13d
je cfar_stats_worker_done
mov edi,r12d
call cfar_record_ptr
lea rbx,[rax+1]
xor ecx,ecx
cfar_stats_worker_pixel:
cmp ecx,1024
je cfar_stats_worker_next
movzx eax,byte [rbx+rcx]
add qword [r14+0],rax
imul eax,eax
add qword [r15+0],rax
movzx eax,byte [rbx+rcx+1024]
add qword [r14+8],rax
imul eax,eax
add qword [r15+8],rax
movzx eax,byte [rbx+rcx+2048]
add qword [r14+16],rax
imul eax,eax
add qword [r15+16],rax
inc ecx
jmp cfar_stats_worker_pixel
cfar_stats_worker_next:
inc r12d
jmp cfar_stats_worker_record
cfar_stats_worker_done:
xor eax,eax
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_zero_optimizer:
xor ebx,ebx
cfar_zero_optimizer_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_zero_optimizer_done
xor r12d,r12d
cfar_zero_optimizer_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_zero_optimizer_next_gpu
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_M]
mov rdx,qword [r13+PD_ELEMS]
call cfar_zero_gpu_f32
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_V]
mov rdx,qword [r13+PD_ELEMS]
call cfar_zero_gpu_f32
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_GRAD]
mov rdx,qword [r13+PD_ELEMS]
call cfar_zero_gpu_f32
inc r12d
jmp cfar_zero_optimizer_param
cfar_zero_optimizer_next_gpu:
inc ebx
jmp cfar_zero_optimizer_gpu
cfar_zero_optimizer_done:
ret
cfar_zero_gpu_f32:
; rdi device,rsi gpu span,rdx elements.
push rbx
push r12
lea rsp,[rsp-96]
mov rbx,rdi
mov r12,rdx
lea rdi,[rsp]
; rsi preserved as incoming span.
xor edx,edx
lea rcx,[r12*4]
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+GA_SIZE]
mov rsi,r12
call gpu_arg_u64
mov rdi,rbx
lea rsi,[rel cfar_k_zero_f32]
mov rdx,rsp
mov rcx,2
mov r8,r12
mov r9,256
call launch_1d
lea rsp,[rsp+96]
pop r12
pop rbx
ret
cfar_train:
call cfar_download
call cfar_runtime_open
call cfar_checkpoint_load_try
cmp eax,0
je cfar_train_resume
cmp eax,1
jne cfar_train_bad
call cfar_init_params
call cfar_zero_optimizer
cfar_train_resume:
mov rax,qword [rel cfar_step]
cmp rax,CFAR_TOTAL_STEPS
jae cfar_train_final
cfar_train_loop:
mov rax,qword [rel cfar_step]
xor edx,edx
mov rcx,CFAR_STEPS_PER_EPOCH
div rcx
mov qword [rel cfar_epoch],rax
mov r12d,edx
call cfar_train_step
cmp eax,1
je cfar_train_loop
test eax,eax
jnz cfar_train_bad
inc qword [rel cfar_step]
mov rax,qword [rel cfar_step]
xor edx,edx
mov rcx,CFAR_CHECKPOINT_PERIOD
div rcx
test rdx,rdx
jnz cfar_train_no_ckpt
call cfar_checkpoint_write
cfar_train_no_ckpt:
mov rax,qword [rel cfar_step]
cmp rax,CFAR_TOTAL_STEPS
jb cfar_train_loop
cfar_train_final:
call cfar_checkpoint_write
xor eax,eax
ret
cfar_train_bad:
lea rdi,[rel cfar_err_train]
mov rsi,cfar_err_train_n
mov edx,70
jmp fatal
cfar_train_step:
; r12d is step within epoch.
call cfar_zero_grads
call cfar_zero_tensor_grads
xor ebx,ebx
cfar_train_step_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_train_step_sync
mov edi,ebx
mov esi,r12d
xor edx,edx
xor ecx,ecx
call cfar_fill_stage
mov edi,ebx
xor esi,esi
call cfar_copy_stage_gpu
mov edi,ebx
xor esi,esi
mov edx,1
call cfar_preprocess
mov edi,ebx
call cfar_forward
mov edi,ebx
xor esi,esi
mov edx,1
call cfar_xent
mov edi,ebx
call cfar_backward
inc ebx
jmp cfar_train_step_gpu
cfar_train_step_sync:
xor ebx,ebx
cfar_train_step_sync_loop:
cmp ebx,CFAR_GPU_MAX
je cfar_train_reduce
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call gpu.sync
test eax,eax
jnz cfar_train_step_bad
inc ebx
jmp cfar_train_step_sync_loop
cfar_train_reduce:
call cfar_reduce_update
cmp eax,1
je cfar_train_step_overflow
test eax,eax
jnz cfar_train_step_bad
xor eax,eax
ret
cfar_train_step_overflow:
mov eax,1
ret
cfar_train_step_bad:
mov eax,-1
ret
cfar_reduce_update:
call cfar_gradients_finite
test eax,eax
jz cfar_reduce_overflow
call cfar_learning_rate
movss dword [rel cfar_lr_current],xmm0
xor r12d,r12d
cfar_reduce_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_reduce_done
xor ebx,ebx
mov edi,0
mov esi,r12d
call cfar_param_desc
mov r13,rax
cfar_reduce_peer:
inc ebx
cmp ebx,CFAR_GPU_MAX
jae cfar_reduce_apply
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r14,rax
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [rel cfar_reduce_scratch]
mov rdx,qword [rel cfar_gpu_handles+rbx*8]
mov rcx,qword [r14+PD_GRAD]
mov r8,qword [r14+PD_ELEMS]
shl r8,2
call gpu.copy_peer
test eax,eax
jnz cfar_reduce_bad
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [r13+PD_GRAD]
mov rdx,qword [rel cfar_reduce_scratch]
mov rcx,qword [r13+PD_ELEMS]
call cfar_launch_add_f32
jmp cfar_reduce_peer
cfar_reduce_apply:
mov rdi,r13
call cfar_launch_adamw
xor ebx,ebx
cfar_reduce_broadcast:
inc ebx
cmp ebx,CFAR_GPU_MAX
jae cfar_reduce_next_param
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r14,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r14+PD_W]
mov rdx,qword [rel cfar_gpu_handles]
mov rcx,qword [r13+PD_W]
mov r8,qword [r13+PD_ELEMS]
shl r8,1
call gpu.copy_peer
test eax,eax
jnz cfar_reduce_bad
jmp cfar_reduce_broadcast
cfar_reduce_next_param:
inc r12d
jmp cfar_reduce_param
cfar_reduce_done:
mov rax,qword [rel cfar_step]
and rax,255
jnz cfar_reduce_done_no_scale
mov rax,qword [rel cfar_loss_scale]
cmp rax,1048576
jae cfar_reduce_done_no_scale
shl rax,1
mov qword [rel cfar_loss_scale],rax
cfar_reduce_done_no_scale:
xor eax,eax
ret
cfar_reduce_overflow:
mov rax,qword [rel cfar_loss_scale]
cmp rax,1
jbe cfar_reduce_bad
shr rax,1
mov qword [rel cfar_loss_scale],rax
call cfar_zero_grads
mov eax,1
ret
cfar_reduce_bad:
mov eax,-1
ret
cfar_gradients_finite:
lea rdi,[rel cfar_finite_flag_host]
mov rsi,CFAR_GPU_MAX*4
call memzero
xor ebx,ebx
cfar_gradients_finite_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_gradients_finite_result
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_finite_flag_gpu+rbx*8]
mov rdx,1
call cfar_zero_gpu_f32
xor r12d,r12d
cfar_gradients_finite_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_gradients_finite_copy
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_GRAD]
mov rdx,qword [r13+PD_ELEMS]
mov rcx,qword [rel cfar_finite_flag_gpu+rbx*8]
call cfar_launch_finite
inc r12d
jmp cfar_gradients_finite_param
cfar_gradients_finite_copy:
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call gpu.sync
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_finite_flag_host+rbx*4]
mov rdx,qword [rel cfar_finite_flag_gpu+rbx*8]
mov rcx,4
call gpu.copy_d2h
inc ebx
jmp cfar_gradients_finite_gpu
cfar_gradients_finite_result:
xor ebx,ebx
cfar_gradients_finite_check:
cmp ebx,CFAR_GPU_MAX
je cfar_gradients_finite_yes
cmp dword [rel cfar_finite_flag_host+rbx*4],0
jne cfar_gradients_finite_no
inc ebx
jmp cfar_gradients_finite_check
cfar_gradients_finite_yes:
mov eax,1
ret
cfar_gradients_finite_no:
xor eax,eax
ret
cfar_launch_add_f32:
; rdi device,rsi dst,rdx src,rcx elements.
push rbx
push r12
push r13
push r14
lea rsp,[rsp-128]
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov r14,rcx
lea rdi,[rsp+0*GA_SIZE]
mov rsi,r12
xor edx,edx
lea rcx,[r14*4]
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,r13
xor edx,edx
lea rcx,[r14*4]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,r14
call gpu_arg_u64
mov rdi,rbx
lea rsi,[rel cfar_k_accum_f32]
mov rdx,rsp
mov rcx,3
mov r8,r14
mov r9,256
call launch_1d
lea rsp,[rsp+128]
pop r14
pop r13
pop r12
pop rbx
ret
cfar_launch_finite:
; rdi device,rsi grad,rdx elements,rcx flag on GPU0 or peer-visible.
push rbx
push r12
push r13
push r14
lea rsp,[rsp-128]
mov rbx,rdi
mov r12,rsi
mov r13,rdx
mov r14,rcx
lea rdi,[rsp]
mov rsi,r12
xor edx,edx
lea rcx,[r13*4]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+GA_SIZE]
mov rsi,r14
xor edx,edx
mov rcx,4
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,r13
call gpu_arg_u64
mov rdi,rbx
lea rsi,[rel cfar_k_finite]
mov rdx,rsp
mov rcx,3
mov r8,r13
mov r9,256
call launch_1d
lea rsp,[rsp+128]
pop r14
pop r13
pop r12
pop rbx
ret
cfar_learning_rate:
mov rax,qword [rel cfar_step]
cmp rax,CFAR_WARMUP_STEPS
jae cfar_learning_rate_cos
cvtsi2ss xmm0,rax
addss xmm0,dword [rel f_one]
mov eax,CFAR_WARMUP_STEPS
cvtsi2ss xmm1,eax
divss xmm0,xmm1
mulss xmm0,dword [rel cfar_lr_max]
ret
cfar_learning_rate_cos:
sub rax,CFAR_WARMUP_STEPS
cvtsi2ss xmm0,rax
mov eax,CFAR_TOTAL_STEPS-CFAR_WARMUP_STEPS
cvtsi2ss xmm1,eax
divss xmm0,xmm1
mulss xmm0,dword [rel f_pi]
call f32_cos
addss xmm0,dword [rel f_one]
mulss xmm0,dword [rel f_half]
movss xmm1,dword [rel cfar_lr_max]
subss xmm1,dword [rel cfar_lr_min]
mulss xmm0,xmm1
addss xmm0,dword [rel cfar_lr_min]
ret
cfar_launch_adamw:
; rdi parameter descriptor on GPU0.
push rbx
push r12
lea rsp,[rsp-512]
mov rbx,rdi
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [rbx+PD_MASTER]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [rbx+PD_GRAD]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [rbx+PD_M]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [rbx+PD_V]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+4*GA_SIZE]
mov rsi,qword [rbx+PD_EMA]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_RW
call gpu_arg_span
lea rdi,[rsp+5*GA_SIZE]
mov rsi,qword [rbx+PD_W]
xor edx,edx
mov rcx,qword [rbx+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+6*GA_SIZE]
mov rsi,qword [rbx+PD_ELEMS]
call gpu_arg_u64
lea rdi,[rsp+7*GA_SIZE]
mov esi,dword [rel cfar_lr_current]
call gpu_arg_f32
lea rdi,[rsp+8*GA_SIZE]
mov esi,dword [rel cfar_beta1]
call gpu_arg_f32
lea rdi,[rsp+9*GA_SIZE]
mov esi,dword [rel cfar_beta2]
call gpu_arg_f32
lea rdi,[rsp+10*GA_SIZE]
mov esi,dword [rel cfar_adam_eps]
call gpu_arg_f32
lea rdi,[rsp+11*GA_SIZE]
mov esi,dword [rel cfar_weight_decay]
call gpu_arg_f32
lea rdi,[rsp+12*GA_SIZE]
mov esi,dword [rel cfar_ema_decay]
call gpu_arg_f32
lea rdi,[rsp+13*GA_SIZE]
mov rsi,qword [rel cfar_step]
inc rsi
call gpu_arg_u64
lea rdi,[rsp+14*GA_SIZE]
mov rsi,qword [rel cfar_loss_scale]
call gpu_arg_u64
lea rdi,[rsp+15*GA_SIZE]
mov rsi,qword [rbx+PD_KIND]
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles]
lea rsi,[rel cfar_k_adamw]
mov rdx,rsp
mov rcx,16
mov r8,qword [rbx+PD_ELEMS]
mov r9,256
call launch_1d
lea rsp,[rsp+512]
pop r12
pop rbx
ret
section .rodata align=16
f_inv_255 dd 0.00392156862745098
f_255 dd 255.0
f_one_4096 dd 0.000244140625
cfar_lr_max dd 0.0025
cfar_lr_min dd 0.00001
cfar_weight_decay dd 0.005
section .bss align=16
cfar_lr_current resd 1
section .text
cfar_checkpoint_payload_alloc:
cmp qword [rel cfar_ck_payload],0
jne cfar_checkpoint_payload_alloc_done
mov rdi,CFAR_PAYLOAD_BYTES
call alloc
mov qword [rel cfar_ck_payload],rax
cfar_checkpoint_payload_alloc_done:
ret
cfar_checkpoint_gather:
call cfar_checkpoint_payload_alloc
mov r15,qword [rel cfar_ck_payload]
xor r12d,r12d
cfar_checkpoint_gather_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_checkpoint_gather_done
mov edi,0
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov r14,qword [r13+PD_ELEMS]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,r15
mov rdx,qword [r13+PD_MASTER]
lea rcx,[r14*4]
call gpu.copy_d2h
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,r15
mov rdx,qword [r13+PD_M]
lea rcx,[r14*4]
call gpu.copy_d2h
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,r15
mov rdx,qword [r13+PD_V]
lea rcx,[r14*4]
call gpu.copy_d2h
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,r15
mov rdx,qword [r13+PD_EMA]
lea rcx,[r14*4]
call gpu.copy_d2h
lea r15,[r15+r14*4]
inc r12d
jmp cfar_checkpoint_gather_param
cfar_checkpoint_gather_done:
ret
cfar_checkpoint_write:
call cfar_checkpoint_gather
lea rdi,[rel cfar_checkpoint_header]
mov rsi,CFAR_CKPT_HEADER
call memzero
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_MAGIC]
lea rsi,[rel cfar_ck_magic16]
mov rdx,16
call memcpy_data
mov dword [rel cfar_checkpoint_header+CFAR_CK_VERSION_O],CFAR_CK_VERSION
mov dword [rel cfar_checkpoint_header+CFAR_CK_HEADER_O],CFAR_CKPT_HEADER
mov qword [rel cfar_checkpoint_header+CFAR_CK_PAYLOAD_O],CFAR_PAYLOAD_BYTES
mov qword [rel cfar_checkpoint_header+CFAR_CK_PARAM_O],CFAR_PARAM_ELEMS
mov rax,qword [rel cfar_step]
mov qword [rel cfar_checkpoint_header+CFAR_CK_STEP_O],rax
mov rax,qword [rel cfar_epoch]
mov qword [rel cfar_checkpoint_header+CFAR_CK_EPOCH_O],rax
mov rax,qword [rel cfar_loss_scale]
mov qword [rel cfar_checkpoint_header+CFAR_CK_LOSS_SCALE_O],rax
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_RNG_O]
lea rsi,[rel cfar_rng_counter]
mov rdx,32
call memcpy_data
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_MEAN_O]
lea rsi,[rel cfar_mean]
mov rdx,12
call memcpy_data
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_INVSTD_O]
lea rsi,[rel cfar_invstd]
mov rdx,12
call memcpy_data
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_ARCH_O]
lea rsi,[rel cfar_arch_sha]
mov rdx,32
call memcpy_data
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_DATA_O]
lea rsi,[rel cfar_dataset_sha]
mov rdx,32
call memcpy_data
call os.clock_ns
mov qword [rel cfar_checkpoint_header+CFAR_CK_CREATED_NS_O],rax
mov rax,qword [rel cfar_ckpt_content_flags]
mov qword [rel cfar_checkpoint_header+CFAR_CK_FLAGS_O],rax
lea rdi,[rel cfar_tmp_sha_ctx]
call sha256_init
lea rdi,[rel cfar_tmp_sha_ctx]
mov rsi,qword [rel cfar_ck_payload]
mov rdx,CFAR_PAYLOAD_BYTES
call sha256_update
lea rdi,[rel cfar_tmp_sha_ctx]
lea rsi,[rel cfar_checkpoint_header+CFAR_CK_PAYLOAD_SHA_O]
call sha256_final
lea rdi,[rel cfar_checkpoint_header+CFAR_CK_HEADER_SHA_O]
mov rsi,32
call memzero
lea rdi,[rel cfar_tmp_sha_ctx]
call sha256_init
lea rdi,[rel cfar_tmp_sha_ctx]
lea rsi,[rel cfar_checkpoint_header]
mov rdx,CFAR_CKPT_HEADER
call sha256_update
lea rdi,[rel cfar_tmp_sha_ctx]
lea rsi,[rel cfar_checkpoint_header+CFAR_CK_HEADER_SHA_O]
call sha256_final
lea rdi,[rel cfar_ckpt_tmp_path]
call os.open_rw_create
test eax,eax
js cfar_checkpoint_write_bad
mov ebx,eax
mov edi,ebx
lea rsi,[rel cfar_checkpoint_header]
mov rdx,CFAR_CKPT_HEADER
call file_write_all
test eax,eax
jnz cfar_checkpoint_write_close_bad
mov edi,ebx
mov rsi,qword [rel cfar_ck_payload]
mov rdx,CFAR_PAYLOAD_BYTES
call file_write_all
test eax,eax
jnz cfar_checkpoint_write_close_bad
mov edi,ebx
call os.fsync
test eax,eax
jnz cfar_checkpoint_write_close_bad
mov edi,ebx
call os.close
lea rdi,[rel cfar_ckpt_tmp_path]
lea rsi,[rel cfar_ckpt_path]
call os.rename
test eax,eax
js cfar_checkpoint_write_bad
lea rdi,[rel cfar_checkpoint_dir]
call os.open_dir
test eax,eax
js cfar_checkpoint_write_bad
mov ebx,eax
mov edi,ebx
call os.fsync
mov edi,ebx
call os.close
xor eax,eax
ret
cfar_checkpoint_write_close_bad:
mov edi,ebx
call os.close
cfar_checkpoint_write_bad:
mov eax,-1
ret
cfar_checkpoint_load_try:
call cfar_checkpoint_payload_alloc
lea rdi,[rel cfar_ckpt_path]
call os.open_ro
test eax,eax
js cfar_checkpoint_load_missing
mov ebx,eax
mov edi,ebx
call os.seek_end
cmp rax,CFAR_CKPT_HEADER+CFAR_PAYLOAD_BYTES
jne cfar_checkpoint_load_bad_close
mov edi,ebx
lea rsi,[rel cfar_checkpoint_header]
mov rdx,CFAR_CKPT_HEADER
xor ecx,ecx
call file_pread_exact
test eax,eax
jnz cfar_checkpoint_load_bad_close
mov edi,ebx
mov rsi,qword [rel cfar_ck_payload]
mov rdx,CFAR_PAYLOAD_BYTES
mov rcx,CFAR_CKPT_HEADER
call file_pread_exact
test eax,eax
jnz cfar_checkpoint_load_bad_close
mov edi,ebx
call os.close
mov rax,qword [rel cfar_checkpoint_header+CFAR_CK_STEP_O]
cmp rax,CFAR_TOTAL_STEPS
ja cfar_checkpoint_load_bad
mov qword [rel cfar_step],rax
mov rax,qword [rel cfar_checkpoint_header+CFAR_CK_EPOCH_O]
mov qword [rel cfar_epoch],rax
mov rax,qword [rel cfar_checkpoint_header+CFAR_CK_LOSS_SCALE_O]
test rax,rax
jz cfar_checkpoint_load_bad
mov qword [rel cfar_loss_scale],rax
lea rdi,[rel cfar_rng_counter]
lea rsi,[rel cfar_checkpoint_header+CFAR_CK_RNG_O]
mov rdx,32
call memcpy_data
lea rdi,[rel cfar_mean]
lea rsi,[rel cfar_checkpoint_header+CFAR_CK_MEAN_O]
mov rdx,12
call memcpy_data
lea rdi,[rel cfar_invstd]
lea rsi,[rel cfar_checkpoint_header+CFAR_CK_INVSTD_O]
mov rdx,12
call memcpy_data
call cfar_checkpoint_scatter
xor eax,eax
ret
cfar_checkpoint_load_missing:
mov eax,1
ret
cfar_checkpoint_load_bad_close:
mov edi,ebx
call os.close
cfar_checkpoint_load_bad:
mov eax,-1
ret
cfar_checkpoint_scatter:
mov r15,qword [rel cfar_ck_payload]
xor r12d,r12d
cfar_checkpoint_scatter_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_checkpoint_scatter_broadcast
mov edi,0
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov r14,qword [r13+PD_ELEMS]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [r13+PD_MASTER]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [r13+PD_M]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [r13+PD_V]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [r13+PD_EMA]
mov rdx,r15
lea rcx,[r14*4]
call gpu.copy_h2d
lea r15,[r15+r14*4]
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,r13
call cfar_cast_param_bf16
inc r12d
jmp cfar_checkpoint_scatter_param
cfar_checkpoint_scatter_broadcast:
xor r12d,r12d
cfar_checkpoint_scatter_bp:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_checkpoint_scatter_done
mov edi,0
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov ebx,1
cfar_checkpoint_scatter_bg:
cmp ebx,CFAR_GPU_MAX
je cfar_checkpoint_scatter_bn
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r14,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r14+PD_W]
mov rdx,qword [rel cfar_gpu_handles]
mov rcx,qword [r13+PD_W]
mov r8,qword [r13+PD_ELEMS]
shl r8,1
call gpu.copy_peer
inc ebx
jmp cfar_checkpoint_scatter_bg
cfar_checkpoint_scatter_bn:
inc r12d
jmp cfar_checkpoint_scatter_bp
cfar_checkpoint_scatter_done:
ret
cfar_load_checkpoint:
call cfar_download
call cfar_runtime_open
call cfar_checkpoint_load_try
test eax,eax
jnz cfar_load_checkpoint_bad
call cfar_use_ema
ret
cfar_load_checkpoint_bad:
lea rdi,[rel cfar_err_ckpt]
mov rsi,cfar_err_ckpt_n
mov edx,66
jmp fatal
cfar_use_ema:
xor r12d,r12d
cfar_use_ema_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_use_ema_done
mov edi,0
mov esi,r12d
call cfar_param_desc
mov r13,rax
; cast EMA into GPU0 BF16 weight.
lea rsp,[rsp-160]
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r13+PD_EMA]
xor edx,edx
mov rcx,qword [r13+PD_ELEMS]
shl rcx,2
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r13+PD_W]
xor edx,edx
mov rcx,qword [r13+PD_ELEMS]
shl rcx,1
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [r13+PD_ELEMS]
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles]
lea rsi,[rel cfar_k_f32_to_bf16]
mov rdx,rsp
mov rcx,3
mov r8,qword [r13+PD_ELEMS]
mov r9,256
call launch_1d
lea rsp,[rsp+160]
mov ebx,1
cfar_use_ema_broadcast:
cmp ebx,CFAR_GPU_MAX
je cfar_use_ema_next
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r14,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r14+PD_W]
mov rdx,qword [rel cfar_gpu_handles]
mov rcx,qword [r13+PD_W]
mov r8,qword [r13+PD_ELEMS]
shl r8,1
call gpu.copy_peer
inc ebx
jmp cfar_use_ema_broadcast
cfar_use_ema_next:
inc r12d
jmp cfar_use_ema_param
cfar_use_ema_done:
ret
cfar_eval:
call cfar_use_ema
mov qword [rel cfar_eval_correct],0
mov qword [rel cfar_eval_total],0
mov qword [rel cfar_eval_nll_f64],0
lea rdi,[rel cfar_eval_confusion]
mov rsi,800
call memzero
xor r12d,r12d
cfar_eval_batch:
cmp r12d,10
je cfar_eval_print
xor ebx,ebx
cfar_eval_gpu_loop:
cmp ebx,CFAR_GPU_MAX
je cfar_eval_sync
mov edi,ebx
mov esi,r12d
xor edx,edx
mov ecx,1
call cfar_fill_stage
mov edi,ebx
xor esi,esi
call cfar_copy_stage_gpu
mov edi,ebx
xor esi,esi
xor edx,edx
call cfar_preprocess
mov edi,ebx
call cfar_forward
mov edi,ebx
xor esi,esi
xor edx,edx
call cfar_xent
mov edi,ebx
call cfar_launch_eval
inc ebx
jmp cfar_eval_gpu_loop
cfar_eval_sync:
xor ebx,ebx
cfar_eval_collect:
cmp ebx,CFAR_GPU_MAX
je cfar_eval_next_batch
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call gpu.sync
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_eval_host+rbx*8]
mov rdx,qword [rel cfar_eval_gpu+rbx*8]
mov rcx,CFAR_LOCAL_BATCH*16
call gpu.copy_d2h
mov rdi,qword [rel cfar_eval_host+rbx*8]
mov rsi,CFAR_VALID_PER_GPU
call cfar_accumulate_eval
inc ebx
jmp cfar_eval_collect
cfar_eval_next_batch:
inc r12d
jmp cfar_eval_batch
cfar_eval_print:
call cfar_print_eval
ret
cfar_launch_eval:
; edi GPU; writes [prediction u32,label u32,loss f32,prob f32].
push rbx
lea rsp,[rsp-192]
mov ebx,edi
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-2
call cfar_tensor_desc
mov r12,rax
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-1
call cfar_tensor_desc
mov r13,rax
lea rdi,[rsp+0*GA_SIZE]
mov rsi,qword [r12+T_CAP]
xor edx,edx
mov rcx,qword [r12+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+1*GA_SIZE]
mov rsi,qword [r13+T_CAP]
xor edx,edx
mov rcx,qword [r13+T_BYTES]
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+2*GA_SIZE]
mov rsi,qword [rel cfar_labels_gpu+rbx*16]
xor edx,edx
mov rcx,CFAR_LOCAL_BATCH
mov r8,GA_SPAN_R
call gpu_arg_span
lea rdi,[rsp+3*GA_SIZE]
mov rsi,qword [rel cfar_eval_gpu+rbx*8]
xor edx,edx
mov rcx,CFAR_LOCAL_BATCH*16
mov r8,GA_SPAN_W
call gpu_arg_span
lea rdi,[rsp+4*GA_SIZE]
mov rsi,CFAR_LOCAL_BATCH
call gpu_arg_u64
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
lea rsi,[rel cfar_k_eval]
mov rdx,rsp
mov rcx,5
mov r8,CFAR_LOCAL_BATCH
mov r9,128
call launch_1d
lea rsp,[rsp+192]
pop rbx
ret
cfar_accumulate_eval:
; rdi records,rsi count.
xor ecx,ecx
cfar_accumulate_eval_loop:
cmp rcx,rsi
je cfar_accumulate_eval_done
mov eax,dword [rdi+rcx*16+0]
mov edx,dword [rdi+rcx*16+4]
cmp edx,10
jae cfar_accumulate_eval_next
inc qword [rel cfar_eval_total]
cmp eax,edx
jne cfar_accumulate_eval_conf
inc qword [rel cfar_eval_correct]
cfar_accumulate_eval_conf:
imul rdx,10
add rdx,rax
inc qword [rel cfar_eval_confusion+rdx*8]
movss xmm0,dword [rdi+rcx*16+8]
cvtss2sd xmm0,xmm0
addsd xmm0,qword [rel cfar_eval_nll_f64]
movsd qword [rel cfar_eval_nll_f64],xmm0
cfar_accumulate_eval_next:
inc rcx
jmp cfar_accumulate_eval_loop
cfar_accumulate_eval_done:
ret
cfar_print_eval:
lea rsp,[rsp-512]
lea rdi,[rsp]
lea rsi,[rel cfar_eval_prefix]
call strlen
mov rdx,rax
call http_append
mov rdi,rax
mov rsi,qword [rel cfar_eval_correct]
call u64_dec
add rdi,rax
lea rsi,[rel cfar_eval_mid]
call strlen
mov rdx,rax
call http_append
mov rdi,rax
mov rax,qword [rel cfar_eval_correct]
cvtsi2sd xmm0,rax
mov rax,qword [rel cfar_eval_total]
cvtsi2sd xmm1,rax
divsd xmm0,xmm1
call f64_fixed6
add rdi,rax
lea rsi,[rel cfar_eval_nll]
call strlen
mov rdx,rax
call http_append
mov rdi,rax
movsd xmm0,qword [rel cfar_eval_nll_f64]
mov rax,qword [rel cfar_eval_total]
cvtsi2sd xmm1,rax
divsd xmm0,xmm1
call f64_fixed6
add rdi,rax
lea rsi,[rel cfar_eval_end]
call strlen
mov rdx,rax
call http_append
mov rdx,rax
sub rdx,rsp
mov edi,1
mov rsi,rsp
call file_write_all
lea rsp,[rsp+512]
ret
f64_fixed6:
; rdi output,xmm0 nonnegative; rax bytes.
push rbx
push r12
mov rbx,rdi
mulsd xmm0,qword [rel f64_million]
addsd xmm0,qword [rel f64_half]
cvttsd2si r12,xmm0
mov rax,r12
xor edx,edx
mov rcx,1000000
div rcx
mov rsi,rax
call u64_dec
add rdi,rax
mov byte [rdi],'.'
inc rdi
mov r12,rdx
mov r8,100000
mov ecx,6
f64_fixed6_loop:
mov rax,r12
xor edx,edx
div r8
add al,'0'
mov byte [rdi],al
inc rdi
mov r12,rdx
mov rax,r8
xor edx,edx
mov r9,10
div r9
mov r8,rax
dec ecx
jnz f64_fixed6_loop
mov rax,rdi
sub rax,rbx
pop r12
pop rbx
ret
section .rodata align=8
f64_million dq 1000000.0
f64_half dq 0.5
section .text
cfar_gpu_lock:
; edi GPU.
lea r8,[rel cfar_gpu_locks+rdi*4]
cfar_gpu_lock_try:
mov eax,0
mov edx,1
lock cmpxchg dword [r8],edx
jz cfar_gpu_lock_ok
mov rdi,r8
mov esi,1
call os.wait32
jmp cfar_gpu_lock_try
cfar_gpu_lock_ok:
ret
cfar_gpu_unlock:
lea r8,[rel cfar_gpu_locks+rdi*4]
mov dword [r8],0
mov rdi,r8
mov esi,1
call os.wake32
ret
cfar_serve:
; rdi port.
push rdi
call os.socket_tcp
test eax,eax
js cfar_serve_bad
mov dword [rel cfar_serve_fd],eax
mov edi,eax
pop rsi
mov edx,1024
call os.bind_listen
test eax,eax
jnz cfar_serve_bad
xor ebx,ebx
cfar_serve_spawn:
cmp ebx,CFAR_SERVE_WORKERS
je cfar_serve_join
mov eax,ebx
shl rax,4
lea r12,[rel cfar_serve_args+rax]
mov eax,dword [rel cfar_serve_fd]
mov dword [r12+0],eax
mov dword [r12+4],ebx
lea rdi,[rel cfar_serve_worker]
mov rsi,r12
call os.thread_start
mov qword [rel cfar_serve_threads+rbx*8],rax
inc ebx
jmp cfar_serve_spawn
cfar_serve_join:
xor ebx,ebx
cfar_serve_join_loop:
cmp ebx,CFAR_SERVE_WORKERS
je cfar_serve_bad
mov rdi,qword [rel cfar_serve_threads+rbx*8]
call os.thread_join
inc ebx
jmp cfar_serve_join_loop
cfar_serve_bad:
mov eax,-1
ret
cfar_serve_worker:
; rdi argument; no hot-path allocation.
push rbx
push r12
push r13
push r14
push r15
mov rbx,rdi
mov r12d,dword [rbx+0]
mov r13d,dword [rbx+4]
mov eax,r13d
imul rax,CFAR_SLOT_SIZE
lea r14,[rel cfar_serve_storage+rax]
mov eax,r13d
and eax,7
mov dword [r14+CFAR_SLOT_GPU],eax
cfar_serve_worker_accept:
mov edi,r12d
call os.accept
test eax,eax
js cfar_serve_worker_accept
mov r15d,eax
mov edi,r15d
lea rsi,[r14+CFAR_SLOT_REQ]
mov rdx,CFAR_SERVE_BUF
lea rcx,[r14+CFAR_SLOT_META]
call http_read_request
test eax,eax
jnz cfar_serve_worker_bad
mov rax,qword [r14+CFAR_SLOT_META+8]
lea rdi,[r14+CFAR_SLOT_REQ+rax]
mov rsi,qword [r14+CFAR_SLOT_META+16]
lea rdx,[rel cfar_health_path]
call cfar_path_equal
cmp eax,1
je cfar_serve_worker_health
mov rax,qword [r14+CFAR_SLOT_META+0]
cmp rax,2
jne cfar_serve_worker_notfound
mov rax,qword [r14+CFAR_SLOT_META+8]
lea rdi,[r14+CFAR_SLOT_REQ+rax]
mov rsi,qword [r14+CFAR_SLOT_META+16]
lea rdx,[rel cfar_predict_path]
call cfar_path_equal
cmp eax,1
jne cfar_serve_worker_notfound
cmp qword [r14+CFAR_SLOT_META+32],CFAR_IMAGE_BYTES
jne cfar_serve_worker_bad
mov rax,qword [r14+CFAR_SLOT_META+24]
lea rsi,[r14+CFAR_SLOT_REQ+rax]
lea rdx,[r14+CFAR_SLOT_JSON]
mov edi,dword [r14+CFAR_SLOT_GPU]
call cfar_serve_one
mov edi,r15d
lea rsi,[rel ct_json]
lea rdx,[r14+CFAR_SLOT_JSON]
mov rcx,rax
call http_send_fixed
jmp cfar_serve_worker_close
cfar_serve_worker_health:
mov edi,r15d
lea rsi,[rel cfar_health_headers]
mov rdx,cfar_health_headers_n
call socket_send_all
mov edi,r15d
lea rsi,[rel cfar_health_json]
mov rdx,cfar_health_json_n
call socket_send_all
jmp cfar_serve_worker_close
cfar_serve_worker_notfound:
mov edi,r15d
lea rsi,[rel s_notfound]
mov rdx,s_notfound_n
call socket_send_all
jmp cfar_serve_worker_close
cfar_serve_worker_bad:
mov edi,r15d
lea rsi,[rel s_bad]
mov rdx,s_bad_n
call socket_send_all
cfar_serve_worker_close:
mov edi,r15d
call os.shutdown
mov edi,r15d
call os.close
jmp cfar_serve_worker_accept
cfar_path_equal:
; rdi request path,rsi bytes,rdx expected z.
push rdi
push rsi
mov rdi,rdx
call strlen
mov rcx,rax
pop rsi
pop rdi
cmp rsi,rcx
jne cfar_path_equal_no
mov rsi,rdx
mov rdx,rcx
call memcmp
test eax,eax
setz al
movzx eax,al
ret
cfar_path_equal_no:
xor eax,eax
ret
cfar_serve_one:
; edi GPU,rsi exact 3072 request bytes,rdx output>=1024; rax JSON bytes.
push rbx
push r12
push r13
push r14
push r15
lea rsp,[rsp-128]
mov ebx,edi
mov r12,rsi
mov r15,rdx
mov edi,ebx
call cfar_gpu_lock
mov eax,ebx
shl eax,1
mov r13d,eax
mov r14,qword [rel cfar_stage_host+r13*8]
mov r13,qword [rel cfar_labels_host+r13*8]
mov rdi,r14
mov rsi,r12
mov rdx,CFAR_IMAGE_BYTES
call memcpy_data
mov byte [r13],0
lea rdi,[r14+CFAR_IMAGE_BYTES]
mov rsi,(CFAR_LOCAL_BATCH-1)*CFAR_IMAGE_BYTES
call memzero
lea rdi,[r13+1]
mov rcx,CFAR_LOCAL_BATCH-1
mov al,255
rep stosb
mov edi,ebx
xor esi,esi
call cfar_copy_stage_gpu
mov edi,ebx
xor esi,esi
xor edx,edx
call cfar_preprocess
mov edi,ebx
call cfar_forward
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call gpu.sync
test eax,eax
jnz cfar_serve_one_fail
mov edi,ebx
mov esi,CFAR_TENSOR_COUNT-2
call cfar_tensor_desc
mov r13,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,rsp
mov rdx,qword [r13+T_CAP]
mov rcx,40
call gpu.copy_d2h
test eax,eax
jnz cfar_serve_one_fail
mov edi,ebx
call cfar_gpu_unlock
mov rdi,rsp
mov rsi,r15
call cfar_logits_json
jmp cfar_serve_one_done
cfar_serve_one_fail:
mov edi,ebx
call cfar_gpu_unlock
xor eax,eax
cfar_serve_one_done:
lea rsp,[rsp+128]
pop r15
pop r14
pop r13
pop r12
pop rbx
ret
cfar_logits_json:
; rdi ten f32 logits,rsi out>=1024; rax bytes.
push rbx
push r12
push r13
push r14
mov rbx,rdi
mov r12,rsi
movss xmm7,dword [rbx]
xor r13d,r13d
mov ecx,1
cfar_logits_max:
cmp ecx,10
je cfar_logits_exp
movss xmm0,dword [rbx+rcx*4]
comiss xmm0,xmm7
jbe cfar_logits_max_next
movss xmm7,xmm0
mov r13d,ecx
cfar_logits_max_next:
inc ecx
jmp cfar_logits_max
cfar_logits_exp:
xorps xmm6,xmm6
xor ecx,ecx
cfar_logits_exp_loop:
movss xmm0,dword [rbx+rcx*4]
subss xmm0,xmm7
call f32_exp
movss dword [rbx+40+rcx*4],xmm0
addss xmm6,xmm0
inc ecx
cmp ecx,10
jb cfar_logits_exp_loop
mov rdi,r12
lea rsi,[rel cfar_json_prefix]
call strlen
mov rdx,rax
call http_append
mov rdi,rax
mov esi,r13d
call u64_dec
add rdi,rax
lea rsi,[rel cfar_json_name]
call strlen
mov rdx,rax
call http_append
mov r14,rax
mov rsi,qword [rel cfar_class_ptrs+r13*8]
mov rdi,rsi
call strlen
mov rdx,rax
mov rdi,r14
call http_append
mov rdi,rax
lea rsi,[rel cfar_json_probs]
call strlen
mov rdx,rax
call http_append
mov r14,rax
xor r13d,r13d
cfar_logits_json_prob:
mov rdi,r14
movss xmm0,dword [rbx+40+r13*4]
divss xmm0,xmm6
cvtss2sd xmm0,xmm0
call f64_fixed6
add r14,rax
inc r13d
cmp r13d,10
je cfar_logits_json_tail
mov byte [r14],','
inc r14
jmp cfar_logits_json_prob
cfar_logits_json_tail:
mov rdi,r14
lea rsi,[rel cfar_json_tail]
call strlen
mov rdx,rax
call http_append
sub rax,r12
pop r14
pop r13
pop r12
pop rbx
ret
cfar_destroy_all:
cmp dword [rel cfar_runtime_live],0
je cfar_destroy_all_done
xor ebx,ebx
cfar_destroy_gpu:
cmp ebx,CFAR_GPU_MAX
je cfar_destroy_global
xor r12d,r12d
cfar_destroy_param:
cmp r12d,CFAR_PARAM_TENSORS
je cfar_destroy_tensor_start
mov edi,ebx
mov esi,r12d
call cfar_param_desc
mov r13,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_W]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_MASTER]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_GRAD]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_M]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_V]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [r13+PD_EMA]
call gpu.free
inc r12d
jmp cfar_destroy_param
cfar_destroy_tensor_start:
xor r12d,r12d
cfar_destroy_tensor:
cmp r12d,CFAR_TENSOR_COUNT
je cfar_destroy_aux_start
mov edi,ebx
mov esi,r12d
call cfar_tensor_desc
mov rsi,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call tensor_gpu_free
mov edi,ebx
mov esi,r12d
call cfar_grad_desc
mov rsi,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call tensor_gpu_free
inc r12d
jmp cfar_destroy_tensor
cfar_destroy_aux_start:
xor r12d,r12d
cfar_destroy_aux:
cmp r12d,CFAR_OP_COUNT
je cfar_destroy_stage
mov edi,ebx
mov esi,r12d
call cfar_aux_desc
mov rsi,rax
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
call tensor_gpu_free
inc r12d
jmp cfar_destroy_aux
cfar_destroy_stage:
xor r12d,r12d
cfar_destroy_stage_parity:
cmp r12d,2
je cfar_destroy_eval
mov eax,ebx
shl eax,1
add eax,r12d
mov r13d,eax
mov rdi,qword [rel cfar_stage_host+r13*8]
call free
mov rdi,qword [rel cfar_labels_host+r13*8]
call free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_stage_gpu+r13*8]
call gpu.free
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_labels_gpu+r13*8]
call gpu.free
inc r12d
jmp cfar_destroy_stage_parity
cfar_destroy_eval:
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_eval_gpu+rbx*8]
call gpu.free
mov rdi,qword [rel cfar_eval_host+rbx*8]
call free
inc ebx
jmp cfar_destroy_gpu
cfar_destroy_global:
mov rdi,qword [rel cfar_gpu_handles]
mov rsi,qword [rel cfar_reduce_scratch]
call gpu.free
xor ebx,ebx
cfar_destroy_flags:
cmp ebx,CFAR_GPU_MAX
je cfar_destroy_flags_done
mov rdi,qword [rel cfar_gpu_handles+rbx*8]
mov rsi,qword [rel cfar_finite_flag_gpu+rbx*8]
call gpu.free
inc ebx
jmp cfar_destroy_flags
cfar_destroy_flags_done:
call cfar_close_gpus
call cfar_unmap_batches
mov dword [rel cfar_runtime_live],0
cfar_destroy_all_done:
ret
; ---- SM89 kernels. Final SASS is accepted only after span-confinement verification. ----
cfar_k_zero_f32:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_zero_f32(
.param .u64 p0,
.param .u64 p1
)
{
.reg .pred %p<2>;
.reg .b32 %r<5>;
.reg .b64 %rd<5>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd2,%r3;
setp.ge.u64 %p0,%rd2,%rd1;
@%p0 bra ZF_DONE;
shl.b64 %rd3,%rd2,2;
add.u64 %rd4,%rd0,%rd3;
st.global.u32 [%rd4],0;
ZF_DONE:
ret;
}
.endgpu
cfar_k_f32_to_bf16:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_f32_to_bf16(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2
)
{
.reg .pred %p<2>;
.reg .b16 %h<2>;
.reg .b32 %r<5>;
.reg .b64 %rd<8>;
.reg .f32 %f<2>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd3,%r3;
setp.ge.u64 %p0,%rd3,%rd2;
@%p0 bra CBF_DONE;
shl.b64 %rd4,%rd3,2;
add.u64 %rd5,%rd0,%rd4;
ld.global.f32 %f0,[%rd5];
cvt.rn.bf16.f32 %h0,%f0;
shl.b64 %rd6,%rd3,1;
add.u64 %rd7,%rd1,%rd6;
st.global.b16 [%rd7],%h0;
CBF_DONE:
ret;
}
.endgpu
cfar_k_accum_f32:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_accum_f32(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2
)
{
.reg .pred %p<2>;
.reg .b32 %r<5>;
.reg .b64 %rd<10>;
.reg .f32 %f<3>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd3,%r3;
setp.ge.u64 %p0,%rd3,%rd2;
@%p0 bra AF_DONE;
shl.b64 %rd4,%rd3,2;
add.u64 %rd5,%rd0,%rd4;
add.u64 %rd6,%rd1,%rd4;
ld.global.f32 %f0,[%rd5];
ld.global.f32 %f1,[%rd6];
add.rn.f32 %f2,%f0,%f1;
st.global.f32 [%rd5],%f2;
AF_DONE:
ret;
}
.endgpu
cfar_k_finite:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_finite(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2
)
{
.reg .pred %p<3>;
.reg .b32 %r<8>;
.reg .b64 %rd<8>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd3,%r3;
setp.ge.u64 %p0,%rd3,%rd2;
@%p0 bra FI_DONE;
shl.b64 %rd4,%rd3,2;
add.u64 %rd5,%rd0,%rd4;
ld.global.u32 %r4,[%rd5];
and.b32 %r5,%r4,0x7f800000;
setp.eq.u32 %p1,%r5,0x7f800000;
@!%p1 bra FI_DONE;
atom.global.or.b32 %r6,[%rd1],1;
FI_DONE:
ret;
}
.endgpu
cfar_k_adamw:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_adamw(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5,
.param .u64 p6,
.param .f32 p7,
.param .f32 p8,
.param .f32 p9,
.param .f32 p10,
.param .f32 p11,
.param .f32 p12,
.param .u64 p13,
.param .u64 p14,
.param .u64 p15
)
{
.reg .pred %p<6>;
.reg .b16 %h<2>;
.reg .b32 %r<12>;
.reg .b64 %rd<24>;
.reg .f32 %f<32>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
ld.param.u64 %rd6,[p6];
ld.param.f32 %f0,[p7];
ld.param.f32 %f1,[p8];
ld.param.f32 %f2,[p9];
ld.param.f32 %f3,[p10];
ld.param.f32 %f4,[p11];
ld.param.f32 %f5,[p12];
ld.param.u64 %rd7,[p13];
ld.param.u64 %rd8,[p14];
ld.param.u64 %rd9,[p15];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd10,%r3;
setp.ge.u64 %p0,%rd10,%rd6;
@%p0 bra AD_DONE;
shl.b64 %rd11,%rd10,2;
add.u64 %rd12,%rd0,%rd11;
add.u64 %rd13,%rd1,%rd11;
add.u64 %rd14,%rd2,%rd11;
add.u64 %rd15,%rd3,%rd11;
add.u64 %rd16,%rd4,%rd11;
ld.global.f32 %f6,[%rd12];
ld.global.f32 %f7,[%rd13];
ld.global.f32 %f8,[%rd14];
ld.global.f32 %f9,[%rd15];
ld.global.f32 %f10,[%rd16];
cvt.rn.f32.u64 %f11,%rd8;
mul.rn.f32 %f11,%f11,1000f;
div.rn.f32 %f7,%f7,%f11;
sub.rn.f32 %f12,1f,%f1;
mul.rn.f32 %f13,%f1,%f8;
fma.rn.f32 %f8,%f12,%f7,%f13;
mul.rn.f32 %f14,%f7,%f7;
sub.rn.f32 %f15,1f,%f2;
mul.rn.f32 %f16,%f2,%f9;
fma.rn.f32 %f9,%f15,%f14,%f16;
lg2.approx.f32 %f17,%f1;
cvt.rn.f32.u64 %f18,%rd7;
mul.rn.f32 %f17,%f17,%f18;
ex2.approx.f32 %f17,%f17;
sub.rn.f32 %f17,1f,%f17;
lg2.approx.f32 %f19,%f2;
mul.rn.f32 %f19,%f19,%f18;
ex2.approx.f32 %f19,%f19;
sub.rn.f32 %f19,1f,%f19;
div.rn.f32 %f20,%f8,%f17;
div.rn.f32 %f21,%f9,%f19;
sqrt.rn.f32 %f21,%f21;
add.rn.f32 %f21,%f21,%f3;
div.rn.f32 %f20,%f20,%f21;
setp.eq.u64 %p4,%rd9,2;
setp.eq.u64 %p5,%rd9,3;
@%p4 mov.f32 %f4,0f;
@%p5 mov.f32 %f4,0f;
setp.eq.u64 %p4,%rd9,5;
@%p4 mov.f32 %f4,0f;
fma.rn.f32 %f20,%f4,%f6,%f20;
fma.rn.f32 %f6,-%f0,%f20,%f6;
sub.rn.f32 %f22,1f,%f5;
mul.rn.f32 %f23,%f5,%f10;
fma.rn.f32 %f10,%f22,%f6,%f23;
st.global.f32 [%rd12],%f6;
st.global.f32 [%rd14],%f8;
st.global.f32 [%rd15],%f9;
st.global.f32 [%rd16],%f10;
st.global.u32 [%rd13],0;
cvt.rn.bf16.f32 %h0,%f6;
shl.b64 %rd17,%rd10,1;
add.u64 %rd18,%rd5,%rd17;
st.global.b16 [%rd18],%h0;
AD_DONE:
ret;
}
.endgpu
cfar_k_preprocess:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_preprocess(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5,
.param .f32 p6,
.param .f32 p7,
.param .f32 p8,
.param .f32 p9,
.param .f32 p10,
.param .f32 p11
)
{
.reg .pred %p<12>;
.reg .b16 %h<2>;
.reg .b32 %r<40>;
.reg .b64 %rd<24>;
.reg .f32 %f<12>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
ld.param.f32 %f0,[p6];
ld.param.f32 %f1,[p7];
ld.param.f32 %f2,[p8];
ld.param.f32 %f3,[p9];
ld.param.f32 %f4,[p10];
ld.param.f32 %f5,[p11];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd6,%r3;
mov.u64 %rd7,3072;
mul.lo.u64 %rd8,%rd2,%rd7;
setp.ge.u64 %p0,%rd6,%rd8;
@%p0 bra PP_DONE;
div.u64 %rd9,%rd6,%rd7;
rem.u64 %rd10,%rd6,%rd7;
cvt.u32.u64 %r4,%rd9;
cvt.u32.u64 %r5,%rd10;
div.u32 %r6,%r5,3;
rem.u32 %r7,%r5,3;
div.u32 %r8,%r6,32;
rem.u32 %r9,%r6,32;
cvt.u32.u64 %r26,%rd3;
cvt.u32.u64 %r27,%rd4;
mov.u32 %r10,%r4;
xor.b32 %r10,%r10,%r26;
mad.lo.u32 %r10,%r10,0x9e3779b9,%r27;
xor.b32 %r11,%r10,0x85ebca6b;
shr.u32 %r12,%r11,16;
xor.b32 %r11,%r11,%r12;
rem.u32 %r13,%r11,9;
shr.u32 %r14,%r11,8;
rem.u32 %r14,%r14,9;
shr.u32 %r15,%r11,20;
and.b32 %r15,%r15,1;
setp.eq.u64 %p1,%rd5,0;
@%p1 mov.u32 %r13,4;
@%p1 mov.u32 %r14,4;
@%p1 mov.u32 %r15,0;
sub.s32 %r16,%r8,4;
add.s32 %r16,%r16,%r14;
sub.s32 %r17,%r9,4;
add.s32 %r17,%r17,%r13;
setp.lt.s32 %p2,%r16,0;
setp.ge.s32 %p3,%r16,32;
or.pred %p4,%p2,%p3;
setp.lt.s32 %p5,%r17,0;
setp.ge.s32 %p6,%r17,32;
or.pred %p7,%p5,%p6;
or.pred %p8,%p4,%p7;
setp.eq.u32 %p9,%r15,1;
@%p9 sub.s32 %r17,31,%r17;
; deterministic 8x8 cutout after crop.
shr.u32 %r18,%r11,4;
rem.u32 %r18,%r18,25;
add.u32 %r18,%r18,4;
shr.u32 %r19,%r11,12;
rem.u32 %r19,%r19,25;
add.u32 %r19,%r19,4;
sub.s32 %r20,%r8,%r18;
abs.s32 %r20,%r20;
sub.s32 %r21,%r9,%r19;
abs.s32 %r21,%r21;
setp.lt.s32 %p10,%r20,4;
setp.lt.s32 %p11,%r21,4;
and.pred %p10,%p10,%p11;
setp.ne.u64 %p11,%rd5,0;
and.pred %p10,%p10,%p11;
or.pred %p8,%p8,%p10;
@%p8 mov.f32 %f6,0f;
@%p8 bra PP_STORE;
mad.lo.u32 %r22,%r16,32,%r17;
mad.lo.u32 %r23,%r7,1024,%r22;
mad.lo.u32 %r24,%r4,3072,%r23;
cvt.u64.u32 %rd11,%r24;
add.u64 %rd12,%rd0,%rd11;
ld.global.u8 %r25,[%rd12];
cvt.rn.f32.u32 %f6,%r25;
mul.rn.f32 %f6,%f6,0f3b808081;
PP_NORMALIZE:
setp.eq.u32 %p2,%r7,0;
@%p2 sub.f32 %f6,%f6,%f0;
@%p2 mul.f32 %f6,%f6,%f3;
setp.eq.u32 %p3,%r7,1;
@%p3 sub.f32 %f6,%f6,%f1;
@%p3 mul.f32 %f6,%f6,%f4;
setp.eq.u32 %p4,%r7,2;
@%p4 sub.f32 %f6,%f6,%f2;
@%p4 mul.f32 %f6,%f6,%f5;
PP_STORE:
cvt.rn.bf16.f32 %h0,%f6;
shl.b64 %rd13,%rd6,1;
add.u64 %rd14,%rd1,%rd13;
st.global.b16 [%rd14],%h0;
PP_DONE:
ret;
}
.endgpu
cfar_k_silu_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_silu_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2
)
{
.reg .pred %p<2>;
.reg .b16 %h<3>;
.reg .b32 %r<5>;
.reg .b64 %rd<10>;
.reg .f32 %f<8>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd3,%r3;
setp.ge.u64 %p0,%rd3,%rd2;
@%p0 bra SF_DONE;
shl.b64 %rd4,%rd3,1;
add.u64 %rd5,%rd0,%rd4;
add.u64 %rd6,%rd1,%rd4;
ld.global.b16 %h0,[%rd5];
cvt.f32.bf16 %f0,%h0;
neg.f32 %f1,%f0;
mul.f32 %f1,%f1,1.4426950408889634f;
ex2.approx.f32 %f1,%f1;
add.f32 %f1,%f1,1f;
rcp.approx.f32 %f1,%f1;
mul.f32 %f2,%f0,%f1;
cvt.rn.bf16.f32 %h1,%f2;
st.global.b16 [%rd6],%h1;
SF_DONE:
ret;
}
.endgpu
cfar_k_add_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_add_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3
)
{
.reg .pred %p<2>;
.reg .b16 %h<4>;
.reg .b32 %r<5>;
.reg .b64 %rd<12>;
.reg .f32 %f<4>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd4,%r3;
setp.ge.u64 %p0,%rd4,%rd3;
@%p0 bra ADF_DONE;
shl.b64 %rd5,%rd4,1;
add.u64 %rd6,%rd0,%rd5;
add.u64 %rd7,%rd1,%rd5;
add.u64 %rd8,%rd2,%rd5;
ld.global.b16 %h0,[%rd6];
ld.global.b16 %h1,[%rd7];
cvt.f32.bf16 %f0,%h0;
cvt.f32.bf16 %f1,%h1;
add.f32 %f2,%f0,%f1;
cvt.rn.bf16.f32 %h2,%f2;
st.global.b16 [%rd8],%h2;
ADF_DONE:
ret;
}
.endgpu
cfar_k_conv_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_conv_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5,
.param .u64 p6,
.param .u64 p7,
.param .u64 p8,
.param .u64 p9,
.param .u64 p10
)
{
.reg .pred %p<12>;
.reg .b16 %h<4>;
.reg .b32 %r<48>;
.reg .b64 %rd<40>;
.reg .f32 %f<8>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
ld.param.u64 %rd6,[p6];
ld.param.u64 %rd7,[p7];
ld.param.u64 %rd8,[p8];
ld.param.u64 %rd9,[p9];
ld.param.u64 %rd10,[p10];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd11,%r3;
add.u64 %rd12,%rd4,%rd10;
add.u64 %rd12,%rd12,%rd10;
sub.u64 %rd12,%rd12,%rd8;
div.u64 %rd13,%rd12,%rd9;
add.u64 %rd13,%rd13,1;
add.u64 %rd14,%rd5,%rd10;
add.u64 %rd14,%rd14,%rd10;
sub.u64 %rd14,%rd14,%rd8;
div.u64 %rd15,%rd14,%rd9;
add.u64 %rd15,%rd15,1;
mul.lo.u64 %rd16,%rd3,%rd13;
mul.lo.u64 %rd16,%rd16,%rd15;
mul.lo.u64 %rd16,%rd16,%rd7;
setp.ge.u64 %p0,%rd11,%rd16;
@%p0 bra CVF_DONE;
rem.u64 %rd17,%rd11,%rd7;
div.u64 %rd18,%rd11,%rd7;
rem.u64 %rd19,%rd18,%rd15;
div.u64 %rd18,%rd18,%rd15;
rem.u64 %rd20,%rd18,%rd13;
div.u64 %rd21,%rd18,%rd13;
mov.f32 %f0,0f;
mov.u64 %rd22,0;
CVF_KY:
setp.ge.u64 %p1,%rd22,%rd8;
@%p1 bra CVF_STORE;
mul.lo.u64 %rd23,%rd20,%rd9;
add.u64 %rd23,%rd23,%rd22;
sub.s64 %rd23,%rd23,%rd10;
setp.lt.s64 %p2,%rd23,0;
setp.ge.s64 %p3,%rd23,%rd4;
or.pred %p4,%p2,%p3;
@%p4 bra CVF_KY_NEXT;
mov.u64 %rd24,0;
CVF_KX:
setp.ge.u64 %p5,%rd24,%rd8;
@%p5 bra CVF_KY_NEXT;
mul.lo.u64 %rd25,%rd19,%rd9;
add.u64 %rd25,%rd25,%rd24;
sub.s64 %rd25,%rd25,%rd10;
setp.lt.s64 %p6,%rd25,0;
setp.ge.s64 %p7,%rd25,%rd5;
or.pred %p8,%p6,%p7;
@%p8 bra CVF_KX_NEXT;
mov.u64 %rd26,0;
mul.lo.u64 %rd27,%rd21,%rd4;
add.u64 %rd27,%rd27,%rd23;
mul.lo.u64 %rd27,%rd27,%rd5;
add.u64 %rd27,%rd27,%rd25;
mul.lo.u64 %rd27,%rd27,%rd6;
shl.b64 %rd28,%rd27,1;
add.u64 %rd29,%rd0,%rd28;
mul.lo.u64 %rd30,%rd17,%rd8;
add.u64 %rd30,%rd30,%rd22;
mul.lo.u64 %rd30,%rd30,%rd8;
add.u64 %rd30,%rd30,%rd24;
mul.lo.u64 %rd30,%rd30,%rd6;
shl.b64 %rd31,%rd30,1;
add.u64 %rd32,%rd1,%rd31;
CVF_IC:
setp.ge.u64 %p9,%rd26,%rd6;
@%p9 bra CVF_KX_NEXT;
ld.global.b16 %h0,[%rd29];
ld.global.b16 %h1,[%rd32];
cvt.f32.bf16 %f1,%h0;
cvt.f32.bf16 %f2,%h1;
fma.rn.f32 %f0,%f1,%f2,%f0;
add.u64 %rd29,%rd29,2;
add.u64 %rd32,%rd32,2;
add.u64 %rd26,%rd26,1;
bra CVF_IC;
CVF_KX_NEXT:
add.u64 %rd24,%rd24,1;
bra CVF_KX;
CVF_KY_NEXT:
add.u64 %rd22,%rd22,1;
bra CVF_KY;
CVF_STORE:
cvt.rn.bf16.f32 %h2,%f0;
shl.b64 %rd33,%rd11,1;
add.u64 %rd34,%rd2,%rd33;
st.global.b16 [%rd34],%h2;
CVF_DONE:
ret;
}
.endgpu
cfar_k_gn_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_gn_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5,
.param .u64 p6,
.param .u64 p7,
.param .u64 p8,
.param .u64 p9
)
{
.reg .pred %p<12>;
.reg .b16 %h<5>;
.reg .b32 %r<32>;
.reg .b64 %rd<32>;
.reg .f32 %f<20>;
.shared .align 4 .b8 gs[16];
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
ld.param.u64 %rd6,[p6];
ld.param.u64 %rd7,[p7];
ld.param.u64 %rd8,[p8];
ld.param.u64 %rd9,[p9];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%tid.x;
mov.u32 %r2,%ntid.x;
mul.lo.u64 %rd10,%rd5,%rd9;
cvt.u64.u32 %rd11,%r0;
setp.ge.u64 %p0,%rd11,%rd10;
@%p0 bra GNF_DONE;
setp.ne.u32 %p1,%r1,0;
@%p1 bra GNF_INIT_DONE;
st.shared.u32 [gs+0],0;
st.shared.u32 [gs+4],0;
GNF_INIT_DONE:
bar.sync 0;
rem.u64 %rd12,%rd11,%rd9;
div.u64 %rd13,%rd11,%rd9;
div.u64 %rd14,%rd8,%rd9;
mul.lo.u64 %rd15,%rd6,%rd7;
mul.lo.u64 %rd16,%rd15,%rd14;
cvt.u64.u32 %rd17,%r1;
mov.f32 %f12,0f;
mov.f32 %f13,0f;
GNF_SUM:
setp.ge.u64 %p2,%rd17,%rd16;
@%p2 bra GNF_SUM_DONE;
div.u64 %rd18,%rd17,%rd14;
rem.u64 %rd19,%rd17,%rd14;
mad.lo.u64 %rd20,%rd12,%rd14,%rd19;
mul.lo.u64 %rd21,%rd13,%rd15;
add.u64 %rd21,%rd21,%rd18;
mul.lo.u64 %rd21,%rd21,%rd8;
add.u64 %rd21,%rd21,%rd20;
shl.b64 %rd22,%rd21,1;
add.u64 %rd23,%rd0,%rd22;
ld.global.b16 %h0,[%rd23];
cvt.f32.bf16 %f0,%h0;
mul.f32 %f1,%f0,%f0;
add.f32 %f12,%f12,%f0;
add.f32 %f13,%f13,%f1;
add.u64 %rd17,%rd17,%rd2;
bra GNF_SUM;
GNF_SUM_DONE:
atom.shared.add.f32 %f2,[gs+0],%f12;
atom.shared.add.f32 %f3,[gs+4],%f13;
bar.sync 0;
ld.shared.f32 %f4,[gs+0];
ld.shared.f32 %f5,[gs+4];
cvt.rn.f32.u64 %f6,%rd16;
div.rn.f32 %f4,%f4,%f6;
div.rn.f32 %f5,%f5,%f6;
fma.rn.f32 %f5,-%f4,%f4,%f5;
add.f32 %f5,%f5,0f3727c5ac;
rsqrt.approx.f32 %f7,%f5;
setp.ne.u32 %p3,%r1,0;
@%p3 bra GNF_AUX_DONE;
mul.lo.u64 %rd24,%rd11,8;
add.u64 %rd25,%rd4,%rd24;
st.global.f32 [%rd25+0],%f4;
st.global.f32 [%rd25+4],%f7;
GNF_AUX_DONE:
bar.sync 0;
cvt.u64.u32 %rd17,%r1;
GNF_WRITE:
setp.ge.u64 %p4,%rd17,%rd16;
@%p4 bra GNF_DONE;
div.u64 %rd18,%rd17,%rd14;
rem.u64 %rd19,%rd17,%rd14;
mad.lo.u64 %rd20,%rd12,%rd14,%rd19;
mul.lo.u64 %rd21,%rd13,%rd15;
add.u64 %rd21,%rd21,%rd18;
mul.lo.u64 %rd21,%rd21,%rd8;
add.u64 %rd21,%rd21,%rd20;
shl.b64 %rd22,%rd21,1;
add.u64 %rd23,%rd0,%rd22;
ld.global.b16 %h0,[%rd23];
cvt.f32.bf16 %f0,%h0;
sub.f32 %f8,%f0,%f4;
mul.f32 %f8,%f8,%f7;
shl.b64 %rd26,%rd20,1;
add.u64 %rd27,%rd1,%rd26;
add.u64 %rd28,%rd2,%rd26;
ld.global.b16 %h1,[%rd27];
ld.global.b16 %h2,[%rd28];
cvt.f32.bf16 %f9,%h1;
cvt.f32.bf16 %f10,%h2;
fma.rn.f32 %f11,%f8,%f9,%f10;
cvt.rn.bf16.f32 %h3,%f11;
add.u64 %rd29,%rd3,%rd22;
st.global.b16 [%rd29],%h3;
add.u64 %rd17,%rd17,%rd2;
bra GNF_WRITE;
GNF_DONE:
ret;
}
.endgpu
cfar_k_gap_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_gap_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5
)
{
.reg .pred %p<4>;
.reg .b16 %h<3>;
.reg .b32 %r<8>;
.reg .b64 %rd<24>;
.reg .f32 %f<6>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd6,%r3;
mul.lo.u64 %rd7,%rd2,%rd5;
setp.ge.u64 %p0,%rd6,%rd7;
@%p0 bra GAPF_DONE;
rem.u64 %rd8,%rd6,%rd5;
div.u64 %rd9,%rd6,%rd5;
mul.lo.u64 %rd10,%rd3,%rd4;
mul.lo.u64 %rd17,%rd9,%rd10;
mul.lo.u64 %rd17,%rd17,%rd5;
add.u64 %rd17,%rd17,%rd8;
shl.b64 %rd18,%rd17,1;
add.u64 %rd14,%rd0,%rd18;
shl.b64 %rd15,%rd5,1;
mov.u64 %rd11,0;
mov.f32 %f0,0f;
GAPF_SUM:
setp.ge.u64 %p1,%rd11,%rd10;
@%p1 bra GAPF_STORE;
ld.global.b16 %h0,[%rd14];
cvt.f32.bf16 %f1,%h0;
add.f32 %f0,%f0,%f1;
add.u64 %rd14,%rd14,%rd15;
add.u64 %rd11,%rd11,1;
bra GAPF_SUM;
GAPF_STORE:
cvt.rn.f32.u64 %f2,%rd10;
div.rn.f32 %f0,%f0,%f2;
cvt.rn.bf16.f32 %h1,%f0;
shl.b64 %rd15,%rd6,1;
add.u64 %rd16,%rd1,%rd15;
st.global.b16 [%rd16],%h1;
GAPF_DONE:
ret;
}
.endgpu
cfar_k_linear_fwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_linear_fwd(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .u64 p5,
.param .u64 p6
)
{
.reg .pred %p<4>;
.reg .b16 %h<5>;
.reg .b32 %r<8>;
.reg .b64 %rd<24>;
.reg .f32 %f<8>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.u64 %rd5,[p5];
ld.param.u64 %rd6,[p6];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd7,%r3;
mul.lo.u64 %rd8,%rd4,%rd6;
setp.ge.u64 %p0,%rd7,%rd8;
@%p0 bra LNF_DONE;
rem.u64 %rd9,%rd7,%rd6;
div.u64 %rd10,%rd7,%rd6;
shl.b64 %rd11,%rd9,1;
add.u64 %rd12,%rd2,%rd11;
ld.global.b16 %h0,[%rd12];
cvt.f32.bf16 %f0,%h0;
mov.u64 %rd13,0;
mul.lo.u64 %rd14,%rd10,%rd5;
shl.b64 %rd15,%rd14,1;
add.u64 %rd16,%rd0,%rd15;
mul.lo.u64 %rd17,%rd9,%rd5;
shl.b64 %rd18,%rd17,1;
add.u64 %rd19,%rd1,%rd18;
LNF_K:
setp.ge.u64 %p1,%rd13,%rd5;
@%p1 bra LNF_STORE;
ld.global.b16 %h1,[%rd16];
ld.global.b16 %h2,[%rd19];
cvt.f32.bf16 %f1,%h1;
cvt.f32.bf16 %f2,%h2;
fma.rn.f32 %f0,%f1,%f2,%f0;
add.u64 %rd16,%rd16,2;
add.u64 %rd19,%rd19,2;
add.u64 %rd13,%rd13,1;
bra LNF_K;
LNF_STORE:
shl.b64 %rd20,%rd7,2;
add.u64 %rd21,%rd3,%rd20;
st.global.f32 [%rd21],%f0;
LNF_DONE:
ret;
}
.endgpu
cfar_k_xent:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_xent(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4,
.param .f32 p5,
.param .u64 p6
)
{
.reg .pred %p<8>;
.reg .b32 %r<16>;
.reg .b64 %rd<24>;
.reg .f32 %f<24>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
ld.param.f32 %f0,[p5];
ld.param.u64 %rd5,[p6];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd6,%r3;
setp.ge.u64 %p0,%rd6,%rd4;
@%p0 bra XE_DONE;
add.u64 %rd7,%rd1,%rd6;
ld.global.u8 %r4,[%rd7];
setp.ge.u32 %p1,%r4,10;
@%p1 bra XE_MASK;
mul.lo.u64 %rd8,%rd6,10;
mov.u32 %r5,0;
mov.f32 %f1,-3.402823466e+38f;
XE_MAX:
setp.ge.u32 %p2,%r5,10;
@%p2 bra XE_SUM;
add.u64 %rd9,%rd8,%r5;
shl.b64 %rd10,%rd9,2;
add.u64 %rd11,%rd0,%rd10;
ld.global.f32 %f2,[%rd11];
max.f32 %f1,%f1,%f2;
add.u32 %r5,%r5,1;
bra XE_MAX;
XE_SUM:
mov.u32 %r5,0;
mov.f32 %f3,0f;
XE_SUM_LOOP:
setp.ge.u32 %p3,%r5,10;
@%p3 bra XE_LOSS;
add.u64 %rd9,%rd8,%r5;
shl.b64 %rd10,%rd9,2;
add.u64 %rd11,%rd0,%rd10;
ld.global.f32 %f2,[%rd11];
sub.f32 %f2,%f2,%f1;
mul.f32 %f2,%f2,1.4426950408889634f;
ex2.approx.f32 %f4,%f2;
add.f32 %f3,%f3,%f4;
add.u32 %r5,%r5,1;
bra XE_SUM_LOOP;
XE_LOSS:
lg2.approx.f32 %f5,%f3;
mul.f32 %f5,%f5,0.6931471805599453f;
mov.u32 %r5,0;
mov.f32 %f6,0f;
sub.f32 %f7,1f,%f0;
mov.f32 %f8,0.1111111111111111f;
mul.f32 %f8,%f8,%f0;
cvt.rn.f32.u64 %f9,%rd5;
XE_GRAD:
setp.ge.u32 %p4,%r5,10;
@%p4 bra XE_STORE_LOSS;
add.u64 %rd9,%rd8,%r5;
shl.b64 %rd10,%rd9,2;
add.u64 %rd11,%rd0,%rd10;
add.u64 %rd12,%rd2,%rd10;
ld.global.f32 %f2,[%rd11];
sub.f32 %f10,%f2,%f1;
mul.f32 %f10,%f10,1.4426950408889634f;
ex2.approx.f32 %f10,%f10;
div.rn.f32 %f10,%f10,%f3;
setp.eq.u32 %p5,%r5,%r4;
selp.f32 %f11,%f7,%f8,%p5;
sub.f32 %f12,%f10,%f11;
mul.f32 %f12,%f12,%f9;
st.global.f32 [%rd12],%f12;
sub.f32 %f13,%f2,%f1;
sub.f32 %f13,%f13,%f5;
neg.f32 %f13,%f13;
mul.f32 %f13,%f13,%f11;
add.f32 %f6,%f6,%f13;
add.u32 %r5,%r5,1;
bra XE_GRAD;
XE_STORE_LOSS:
shl.b64 %rd13,%rd6,2;
add.u64 %rd14,%rd3,%rd13;
st.global.f32 [%rd14],%f6;
bra XE_DONE;
XE_MASK:
mul.lo.u64 %rd8,%rd6,10;
mov.u32 %r5,0;
XE_MASK_LOOP:
setp.ge.u32 %p6,%r5,10;
@%p6 bra XE_MASK_LOSS;
add.u64 %rd9,%rd8,%r5;
shl.b64 %rd10,%rd9,2;
add.u64 %rd12,%rd2,%rd10;
st.global.u32 [%rd12],0;
add.u32 %r5,%r5,1;
bra XE_MASK_LOOP;
XE_MASK_LOSS:
shl.b64 %rd13,%rd6,2;
add.u64 %rd14,%rd3,%rd13;
st.global.u32 [%rd14],0;
XE_DONE:
ret;
}
.endgpu
cfar_k_eval:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_eval(
.param .u64 p0,
.param .u64 p1,
.param .u64 p2,
.param .u64 p3,
.param .u64 p4
)
{
.reg .pred %p<5>;
.reg .b32 %r<12>;
.reg .b64 %rd<20>;
.reg .f32 %f<12>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p1];
ld.param.u64 %rd2,[p2];
ld.param.u64 %rd3,[p3];
ld.param.u64 %rd4,[p4];
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd5,%r3;
setp.ge.u64 %p0,%rd5,%rd4;
@%p0 bra EV_DONE;
add.u64 %rd6,%rd2,%rd5;
ld.global.u8 %r4,[%rd6];
mul.lo.u64 %rd7,%rd5,10;
mov.u32 %r5,0;
mov.u32 %r6,0;
mov.f32 %f0,-3.402823466e+38f;
mov.f32 %f1,0f;
EV_MAX:
setp.ge.u32 %p1,%r5,10;
@%p1 bra EV_SUM;
add.u64 %rd8,%rd7,%r5;
shl.b64 %rd9,%rd8,2;
add.u64 %rd10,%rd0,%rd9;
ld.global.f32 %f2,[%rd10];
setp.gt.f32 %p2,%f2,%f0;
@%p2 mov.f32 %f0,%f2;
@%p2 mov.u32 %r6,%r5;
add.u32 %r5,%r5,1;
bra EV_MAX;
EV_SUM:
mov.u32 %r5,0;
EV_SUM_LOOP:
setp.ge.u32 %p3,%r5,10;
@%p3 bra EV_STORE;
add.u64 %rd8,%rd7,%r5;
shl.b64 %rd9,%rd8,2;
add.u64 %rd10,%rd0,%rd9;
ld.global.f32 %f2,[%rd10];
sub.f32 %f2,%f2,%f0;
mul.f32 %f2,%f2,1.4426950408889634f;
ex2.approx.f32 %f2,%f2;
add.f32 %f1,%f1,%f2;
add.u32 %r5,%r5,1;
bra EV_SUM_LOOP;
EV_STORE:
cvt.u64.u32 %rd11,%r4;
add.u64 %rd11,%rd7,%rd11;
shl.b64 %rd11,%rd11,2;
add.u64 %rd12,%rd0,%rd11;
ld.global.f32 %f3,[%rd12];
lg2.approx.f32 %f4,%f1;
mul.f32 %f4,%f4,0.6931471805599453f;
sub.f32 %f5,%f0,%f3;
add.f32 %f3,%f5,%f4;
mul.lo.u64 %rd13,%rd5,16;
add.u64 %rd14,%rd3,%rd13;
st.global.u32 [%rd14+0],%r6;
st.global.u32 [%rd14+4],%r4;
st.global.f32 [%rd14+8],%f3;
st.global.u32 [%rd14+12],0;
EV_DONE:
ret;
}
.endgpu
cfar_k_silu_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.const .align 4 .u32 cfar_op_meta[624] = {
1,32,32,3,32,32,64,3,1,1,0,0,
2,32,32,64,32,32,64,0,0,0,8,1,
3,32,32,64,32,32,64,0,0,0,0,0,
1,32,32,64,32,32,64,3,1,1,0,3,
2,32,32,64,32,32,64,0,0,0,8,4,
3,32,32,64,32,32,64,0,0,0,0,0,
1,32,32,64,32,32,64,3,1,1,0,6,
2,32,32,64,32,32,64,0,0,0,8,7,
4,32,32,64,32,32,64,0,0,0,0,0,
3,32,32,64,32,32,64,0,0,0,0,0,
1,32,32,64,32,32,64,3,1,1,0,9,
2,32,32,64,32,32,64,0,0,0,8,10,
3,32,32,64,32,32,64,0,0,0,0,0,
1,32,32,64,32,32,64,3,1,1,0,12,
2,32,32,64,32,32,64,0,0,0,8,13,
4,32,32,64,32,32,64,0,0,0,0,0,
3,32,32,64,32,32,64,0,0,0,0,0,
1,32,32,64,16,16,128,3,2,1,0,15,
2,16,16,128,16,16,128,0,0,0,8,16,
3,16,16,128,16,16,128,0,0,0,0,0,
1,16,16,128,16,16,128,3,1,1,0,18,
2,16,16,128,16,16,128,0,0,0,8,19,
1,32,32,64,16,16,128,1,2,0,0,21,
2,16,16,128,16,16,128,0,0,0,8,22,
4,16,16,128,16,16,128,0,0,0,0,0,
3,16,16,128,16,16,128,0,0,0,0,0,
1,16,16,128,16,16,128,3,1,1,0,24,
2,16,16,128,16,16,128,0,0,0,8,25,
3,16,16,128,16,16,128,0,0,0,0,0,
1,16,16,128,16,16,128,3,1,1,0,27,
2,16,16,128,16,16,128,0,0,0,8,28,
4,16,16,128,16,16,128,0,0,0,0,0,
3,16,16,128,16,16,128,0,0,0,0,0,
1,16,16,128,8,8,256,3,2,1,0,30,
2,8,8,256,8,8,256,0,0,0,8,31,
3,8,8,256,8,8,256,0,0,0,0,0,
1,8,8,256,8,8,256,3,1,1,0,33,
2,8,8,256,8,8,256,0,0,0,8,34,
1,16,16,128,8,8,256,1,2,0,0,36,
2,8,8,256,8,8,256,0,0,0,8,37,
4,8,8,256,8,8,256,0,0,0,0,0,
3,8,8,256,8,8,256,0,0,0,0,0,
1,8,8,256,8,8,256,3,1,1,0,39,
2,8,8,256,8,8,256,0,0,0,8,40,
3,8,8,256,8,8,256,0,0,0,0,0,
1,8,8,256,8,8,256,3,1,1,0,42,
2,8,8,256,8,8,256,0,0,0,8,43,
4,8,8,256,8,8,256,0,0,0,0,0,
3,8,8,256,8,8,256,0,0,0,0,0,
5,8,8,256,1,1,256,0,0,0,0,0,
6,1,1,256,1,1,10,0,0,0,0,45,
7,1,1,10,1,1,1,0,0,0,0,0
};
.visible .entry cfar_k_silu_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<4>;
.reg .b16 %h<2>;
.reg .b32 %r<12>;
.reg .b64 %rd<20>;
.reg .f32 %f<12>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p2];
ld.param.u64 %rd2,[p3];
ld.param.u64 %rd3,[p11];
ld.param.u64 %rd4,[p14];
mul.lo.u64 %rd5,%rd4,48;
mov.u64 %rd6,cfar_op_meta;
add.u64 %rd6,%rd6,%rd5;
ld.const.u32 %r4,[%rd6+4];
ld.const.u32 %r5,[%rd6+8];
ld.const.u32 %r6,[%rd6+12];
cvt.u64.u32 %rd7,%r4;
mul.lo.u64 %rd7,%rd7,%r5;
mul.lo.u64 %rd7,%rd7,%r6;
mul.lo.u64 %rd7,%rd7,%rd3;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd8,%r3;
setp.ge.u64 %p0,%rd8,%rd7;
@%p0 bra SB_DONE;
shl.b64 %rd9,%rd8,1;
add.u64 %rd10,%rd0,%rd9;
shl.b64 %rd11,%rd8,2;
add.u64 %rd12,%rd1,%rd11;
add.u64 %rd13,%rd2,%rd11;
ld.global.b16 %h0,[%rd10];
cvt.f32.bf16 %f0,%h0;
ld.global.f32 %f1,[%rd12];
neg.f32 %f2,%f0;
mul.f32 %f2,%f2,1.4426950408889634f;
ex2.approx.f32 %f2,%f2;
add.f32 %f2,%f2,1f;
rcp.approx.f32 %f2,%f2;
sub.f32 %f3,1f,%f2;
fma.rn.f32 %f4,%f0,%f3,1f;
mul.f32 %f4,%f4,%f2;
mul.f32 %f5,%f1,%f4;
atom.global.add.f32 %f6,[%rd13],%f5;
SB_DONE:
ret;
}
.endgpu
cfar_k_add_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_add_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<2>;
.reg .b32 %r<10>;
.reg .b64 %rd<20>;
.reg .f32 %f<4>;
ld.param.u64 %rd0,[p2];
ld.param.u64 %rd1,[p3];
ld.param.u64 %rd2,[p5];
ld.param.u64 %rd3,[p11];
ld.param.u64 %rd4,[p14];
mul.lo.u64 %rd5,%rd4,48;
mov.u64 %rd6,cfar_op_meta;
add.u64 %rd6,%rd6,%rd5;
ld.const.u32 %r4,[%rd6+16];
ld.const.u32 %r5,[%rd6+20];
ld.const.u32 %r6,[%rd6+24];
cvt.u64.u32 %rd7,%r4;
mul.lo.u64 %rd7,%rd7,%r5;
mul.lo.u64 %rd7,%rd7,%r6;
mul.lo.u64 %rd7,%rd7,%rd3;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd8,%r3;
setp.ge.u64 %p0,%rd8,%rd7;
@%p0 bra AB_DONE;
shl.b64 %rd9,%rd8,2;
add.u64 %rd10,%rd0,%rd9;
add.u64 %rd11,%rd1,%rd9;
add.u64 %rd12,%rd2,%rd9;
ld.global.f32 %f0,[%rd10];
atom.global.add.f32 %f1,[%rd11],%f0;
atom.global.add.f32 %f2,[%rd12],%f0;
AB_DONE:
ret;
}
.endgpu
cfar_k_gap_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_gap_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<3>;
.reg .b32 %r<12>;
.reg .b64 %rd<24>;
.reg .f32 %f<6>;
ld.param.u64 %rd0,[p2];
ld.param.u64 %rd1,[p3];
ld.param.u64 %rd2,[p11];
ld.param.u64 %rd3,[p14];
mul.lo.u64 %rd4,%rd3,48;
mov.u64 %rd5,cfar_op_meta;
add.u64 %rd5,%rd5,%rd4;
ld.const.u32 %r4,[%rd5+4];
ld.const.u32 %r5,[%rd5+8];
ld.const.u32 %r6,[%rd5+12];
cvt.u64.u32 %rd6,%r4;
mul.lo.u64 %rd6,%rd6,%r5;
mul.lo.u64 %rd7,%rd6,%r6;
mul.lo.u64 %rd7,%rd7,%rd2;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd8,%r3;
setp.ge.u64 %p0,%rd8,%rd7;
@%p0 bra GB_DONE;
rem.u64 %rd9,%rd8,%r6;
div.u64 %rd10,%rd8,%r6;
div.u64 %rd11,%rd10,%rd6;
mul.lo.u64 %rd12,%rd11,%r6;
add.u64 %rd12,%rd12,%rd9;
shl.b64 %rd13,%rd12,2;
add.u64 %rd14,%rd0,%rd13;
ld.global.f32 %f0,[%rd14];
cvt.rn.f32.u64 %f1,%rd6;
div.rn.f32 %f0,%f0,%f1;
shl.b64 %rd15,%rd8,2;
add.u64 %rd16,%rd1,%rd15;
atom.global.add.f32 %f2,[%rd16],%f0;
GB_DONE:
ret;
}
.endgpu
cfar_k_linear_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_linear_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<10>;
.reg .b16 %h<5>;
.reg .b32 %r<20>;
.reg .b64 %rd<40>;
.reg .f32 %f<12>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p2];
ld.param.u64 %rd2,[p3];
ld.param.u64 %rd3,[p6];
ld.param.u64 %rd4,[p7];
ld.param.u64 %rd5,[p9];
ld.param.u64 %rd6,[p11];
ld.param.u64 %rd7,[p14];
mul.lo.u64 %rd8,%rd7,48;
mov.u64 %rd9,cfar_op_meta;
add.u64 %rd9,%rd9,%rd8;
ld.const.u32 %r4,[%rd9+12];
ld.const.u32 %r5,[%rd9+24];
cvt.u64.u32 %rd10,%r4;
cvt.u64.u32 %rd11,%r5;
mul.lo.u64 %rd12,%rd6,%rd10;
mul.lo.u64 %rd13,%rd11,%rd10;
add.u64 %rd14,%rd12,%rd13;
add.u64 %rd15,%rd14,%rd11;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd16,%r3;
setp.ge.u64 %p0,%rd16,%rd15;
@%p0 bra LB_DONE;
setp.lt.u64 %p1,%rd16,%rd12;
@!%p1 bra LB_WEIGHT;
rem.u64 %rd17,%rd16,%rd10;
div.u64 %rd18,%rd16,%rd10;
mul.lo.u64 %rd30,%rd18,%rd11;
shl.b64 %rd31,%rd30,2;
add.u64 %rd22,%rd1,%rd31;
shl.b64 %rd32,%rd17,1;
add.u64 %rd25,%rd3,%rd32;
shl.b64 %rd33,%rd10,1;
mov.u64 %rd19,0;
mov.f32 %f0,0f;
LB_DX_LOOP:
setp.ge.u64 %p2,%rd19,%rd11;
@%p2 bra LB_DX_STORE;
ld.global.f32 %f1,[%rd22];
ld.global.b16 %h0,[%rd25];
cvt.f32.bf16 %f2,%h0;
fma.rn.f32 %f0,%f1,%f2,%f0;
add.u64 %rd22,%rd22,4;
add.u64 %rd25,%rd25,%rd33;
add.u64 %rd19,%rd19,1;
bra LB_DX_LOOP;
LB_DX_STORE:
shl.b64 %rd26,%rd16,2;
add.u64 %rd27,%rd2,%rd26;
atom.global.add.f32 %f3,[%rd27],%f0;
bra LB_DONE;
LB_WEIGHT:
sub.u64 %rd28,%rd16,%rd12;
setp.lt.u64 %p3,%rd28,%rd13;
@!%p3 bra LB_BIAS;
rem.u64 %rd17,%rd28,%rd10;
div.u64 %rd19,%rd28,%rd10;
shl.b64 %rd30,%rd17,1;
add.u64 %rd22,%rd0,%rd30;
shl.b64 %rd31,%rd10,1;
shl.b64 %rd32,%rd19,2;
add.u64 %rd25,%rd1,%rd32;
shl.b64 %rd33,%rd11,2;
mov.u64 %rd18,0;
mov.f32 %f0,0f;
LB_DW_LOOP:
setp.ge.u64 %p4,%rd18,%rd6;
@%p4 bra LB_DW_STORE;
ld.global.b16 %h0,[%rd22];
cvt.f32.bf16 %f1,%h0;
ld.global.f32 %f2,[%rd25];
fma.rn.f32 %f0,%f1,%f2,%f0;
add.u64 %rd22,%rd22,%rd31;
add.u64 %rd25,%rd25,%rd33;
add.u64 %rd18,%rd18,1;
bra LB_DW_LOOP;
LB_DW_STORE:
shl.b64 %rd26,%rd28,2;
add.u64 %rd27,%rd4,%rd26;
st.global.f32 [%rd27],%f0;
bra LB_DONE;
LB_BIAS:
sub.u64 %rd19,%rd28,%rd13;
setp.ge.u64 %p5,%rd19,%rd11;
@%p5 bra LB_DONE;
shl.b64 %rd30,%rd19,2;
add.u64 %rd22,%rd1,%rd30;
shl.b64 %rd31,%rd11,2;
mov.u64 %rd18,0;
mov.f32 %f0,0f;
LB_DB_LOOP:
setp.ge.u64 %p6,%rd18,%rd6;
@%p6 bra LB_DB_STORE;
ld.global.f32 %f1,[%rd22];
add.f32 %f0,%f0,%f1;
add.u64 %rd22,%rd22,%rd31;
add.u64 %rd18,%rd18,1;
bra LB_DB_LOOP;
LB_DB_STORE:
shl.b64 %rd23,%rd19,2;
add.u64 %rd24,%rd5,%rd23;
st.global.f32 [%rd24],%f0;
LB_DONE:
ret;
}
.endgpu
cfar_k_conv_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_conv_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<24>;
.reg .b16 %h<6>;
.reg .b32 %r<32>;
.reg .b64 %rd<64>;
.reg .f32 %f<12>;
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p2];
ld.param.u64 %rd2,[p3];
ld.param.u64 %rd3,[p6];
ld.param.u64 %rd4,[p7];
ld.param.u64 %rd5,[p11];
ld.param.u64 %rd6,[p14];
mul.lo.u64 %rd7,%rd6,48;
mov.u64 %rd8,cfar_op_meta;
add.u64 %rd8,%rd8,%rd7;
ld.const.u32 %r4,[%rd8+4];
ld.const.u32 %r5,[%rd8+8];
ld.const.u32 %r6,[%rd8+12];
ld.const.u32 %r7,[%rd8+16];
ld.const.u32 %r8,[%rd8+20];
ld.const.u32 %r9,[%rd8+24];
ld.const.u32 %r10,[%rd8+28];
ld.const.u32 %r11,[%rd8+32];
ld.const.u32 %r12,[%rd8+36];
cvt.u64.u32 %rd9,%r4;
cvt.u64.u32 %rd10,%r5;
cvt.u64.u32 %rd11,%r6;
cvt.u64.u32 %rd12,%r7;
cvt.u64.u32 %rd13,%r8;
cvt.u64.u32 %rd14,%r9;
cvt.u64.u32 %rd15,%r10;
cvt.u64.u32 %rd16,%r11;
cvt.u64.u32 %rd17,%r12;
mul.lo.u64 %rd18,%rd5,%rd9;
mul.lo.u64 %rd18,%rd18,%rd10;
mul.lo.u64 %rd18,%rd18,%rd11;
mul.lo.u64 %rd19,%rd14,%rd15;
mul.lo.u64 %rd19,%rd19,%rd15;
mul.lo.u64 %rd19,%rd19,%rd11;
add.u64 %rd20,%rd18,%rd19;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%ntid.x;
mov.u32 %r2,%tid.x;
mad.lo.u32 %r3,%r0,%r1,%r2;
cvt.u64.u32 %rd21,%r3;
setp.ge.u64 %p0,%rd21,%rd20;
@%p0 bra CB_DONE;
setp.lt.u64 %p1,%rd21,%rd18;
@!%p1 bra CB_WEIGHT;
rem.u64 %rd22,%rd21,%rd11;
div.u64 %rd23,%rd21,%rd11;
rem.u64 %rd24,%rd23,%rd10;
div.u64 %rd23,%rd23,%rd10;
rem.u64 %rd25,%rd23,%rd9;
div.u64 %rd26,%rd23,%rd9;
mov.f32 %f0,0f;
mov.u64 %rd27,0;
CB_DX_OC:
setp.ge.u64 %p2,%rd27,%rd14;
@%p2 bra CB_DX_STORE;
mov.u64 %rd28,0;
CB_DX_KY:
setp.ge.u64 %p3,%rd28,%rd15;
@%p3 bra CB_DX_OC_NEXT;
add.u64 %rd29,%rd25,%rd17;
setp.lt.u64 %p4,%rd29,%rd28;
@%p4 bra CB_DX_KY_NEXT;
sub.u64 %rd29,%rd29,%rd28;
rem.u64 %rd30,%rd29,%rd16;
setp.ne.u64 %p5,%rd30,0;
@%p5 bra CB_DX_KY_NEXT;
div.u64 %rd31,%rd29,%rd16;
setp.ge.u64 %p6,%rd31,%rd12;
@%p6 bra CB_DX_KY_NEXT;
mov.u64 %rd32,0;
CB_DX_KX:
setp.ge.u64 %p7,%rd32,%rd15;
@%p7 bra CB_DX_KY_NEXT;
add.u64 %rd33,%rd24,%rd17;
setp.lt.u64 %p8,%rd33,%rd32;
@%p8 bra CB_DX_KX_NEXT;
sub.u64 %rd33,%rd33,%rd32;
rem.u64 %rd34,%rd33,%rd16;
setp.ne.u64 %p9,%rd34,0;
@%p9 bra CB_DX_KX_NEXT;
div.u64 %rd35,%rd33,%rd16;
setp.ge.u64 %p10,%rd35,%rd13;
@%p10 bra CB_DX_KX_NEXT;
mul.lo.u64 %rd36,%rd26,%rd12;
add.u64 %rd36,%rd36,%rd31;
mul.lo.u64 %rd36,%rd36,%rd13;
add.u64 %rd36,%rd36,%rd35;
mul.lo.u64 %rd36,%rd36,%rd14;
add.u64 %rd36,%rd36,%rd27;
shl.b64 %rd37,%rd36,2;
add.u64 %rd38,%rd1,%rd37;
ld.global.f32 %f1,[%rd38];
mul.lo.u64 %rd39,%rd27,%rd15;
add.u64 %rd39,%rd39,%rd28;
mul.lo.u64 %rd39,%rd39,%rd15;
add.u64 %rd39,%rd39,%rd32;
mul.lo.u64 %rd39,%rd39,%rd11;
add.u64 %rd39,%rd39,%rd22;
shl.b64 %rd40,%rd39,1;
add.u64 %rd41,%rd3,%rd40;
ld.global.b16 %h0,[%rd41];
cvt.f32.bf16 %f2,%h0;
fma.rn.f32 %f0,%f1,%f2,%f0;
CB_DX_KX_NEXT:
add.u64 %rd32,%rd32,1;
bra CB_DX_KX;
CB_DX_KY_NEXT:
add.u64 %rd28,%rd28,1;
bra CB_DX_KY;
CB_DX_OC_NEXT:
add.u64 %rd27,%rd27,1;
bra CB_DX_OC;
CB_DX_STORE:
shl.b64 %rd42,%rd21,2;
add.u64 %rd43,%rd2,%rd42;
atom.global.add.f32 %f3,[%rd43],%f0;
bra CB_DONE;
CB_WEIGHT:
sub.u64 %rd44,%rd21,%rd18;
setp.ge.u64 %p11,%rd44,%rd19;
@%p11 bra CB_DONE;
rem.u64 %rd22,%rd44,%rd11;
div.u64 %rd45,%rd44,%rd11;
rem.u64 %rd32,%rd45,%rd15;
div.u64 %rd45,%rd45,%rd15;
rem.u64 %rd28,%rd45,%rd15;
div.u64 %rd27,%rd45,%rd15;
mov.u64 %rd26,0;
mov.f32 %f0,0f;
CB_DW_N:
setp.ge.u64 %p12,%rd26,%rd5;
@%p12 bra CB_DW_STORE;
mov.u64 %rd31,0;
CB_DW_OY:
setp.ge.u64 %p13,%rd31,%rd12;
@%p13 bra CB_DW_N_NEXT;
mul.lo.u64 %rd29,%rd31,%rd16;
add.u64 %rd29,%rd29,%rd28;
setp.lt.u64 %p14,%rd29,%rd17;
@%p14 bra CB_DW_OY_NEXT;
sub.u64 %rd29,%rd29,%rd17;
setp.ge.u64 %p15,%rd29,%rd9;
@%p15 bra CB_DW_OY_NEXT;
mul.lo.u64 %rd45,%rd26,%rd9;
add.u64 %rd45,%rd45,%rd29;
mul.lo.u64 %rd45,%rd45,%rd10;
mul.lo.u64 %rd45,%rd45,%rd11;
add.u64 %rd45,%rd45,%rd22;
shl.b64 %rd46,%rd45,1;
add.u64 %rd46,%rd0,%rd46;
mul.lo.u64 %rd45,%rd26,%rd12;
add.u64 %rd45,%rd45,%rd31;
mul.lo.u64 %rd45,%rd45,%rd13;
mul.lo.u64 %rd45,%rd45,%rd14;
add.u64 %rd45,%rd45,%rd27;
shl.b64 %rd41,%rd45,2;
add.u64 %rd41,%rd1,%rd41;
shl.b64 %rd47,%rd11,1;
shl.b64 %rd48,%rd14,2;
sub.u64 %rd33,%rd32,%rd17;
mov.u64 %rd35,0;
CB_DW_OX:
setp.ge.u64 %p16,%rd35,%rd13;
@%p16 bra CB_DW_OY_NEXT;
setp.lt.s64 %p17,%rd33,0;
@%p17 bra CB_DW_OX_NEXT;
setp.ge.u64 %p18,%rd33,%rd10;
@%p18 bra CB_DW_OX_NEXT;
mul.lo.u64 %rd38,%rd33,%rd47;
add.u64 %rd38,%rd46,%rd38;
ld.global.b16 %h0,[%rd38];
cvt.f32.bf16 %f1,%h0;
ld.global.f32 %f2,[%rd41];
fma.rn.f32 %f0,%f1,%f2,%f0;
CB_DW_OX_NEXT:
add.u64 %rd33,%rd33,%rd16;
add.u64 %rd41,%rd41,%rd48;
add.u64 %rd35,%rd35,1;
bra CB_DW_OX;
CB_DW_OY_NEXT:
add.u64 %rd31,%rd31,1;
bra CB_DW_OY;
CB_DW_N_NEXT:
add.u64 %rd26,%rd26,1;
bra CB_DW_N;
CB_DW_STORE:
shl.b64 %rd42,%rd44,2;
add.u64 %rd43,%rd4,%rd42;
st.global.f32 [%rd43],%f0;
CB_DONE:
ret;
}
.endgpu
cfar_k_gn_bwd:
.sm89
.version 8.0
.target sm_89
.address_size 64
.visible .entry cfar_k_gn_bwd(
.param .u64 p0,.param .u64 p1,.param .u64 p2,.param .u64 p3,
.param .u64 p4,.param .u64 p5,.param .u64 p6,.param .u64 p7,
.param .u64 p8,.param .u64 p9,.param .u64 p10,.param .u64 p11,
.param .u64 p12,.param .u64 p13,.param .u64 p14,.param .u64 p15
)
{
.reg .pred %p<16>;
.reg .b16 %h<4>;
.reg .b32 %r<24>;
.reg .b64 %rd<48>;
.reg .f32 %f<24>;
.shared .align 4 .b8 ns[16];
ld.param.u64 %rd0,[p0];
ld.param.u64 %rd1,[p2];
ld.param.u64 %rd2,[p3];
ld.param.u64 %rd3,[p6];
ld.param.u64 %rd4,[p7];
ld.param.u64 %rd5,[p9];
ld.param.u64 %rd6,[p10];
ld.param.u64 %rd7,[p11];
ld.param.u64 %rd8,[p14];
mul.lo.u64 %rd9,%rd8,48;
mov.u64 %rd10,cfar_op_meta;
add.u64 %rd10,%rd10,%rd9;
ld.const.u32 %r4,[%rd10+4];
ld.const.u32 %r5,[%rd10+8];
ld.const.u32 %r6,[%rd10+12];
ld.const.u32 %r7,[%rd10+40];
cvt.u64.u32 %rd11,%r4;
cvt.u64.u32 %rd12,%r5;
cvt.u64.u32 %rd13,%r6;
cvt.u64.u32 %rd14,%r7;
mov.u32 %r0,%ctaid.x;
mov.u32 %r1,%tid.x;
mov.u32 %r2,%ntid.x;
cvt.u64.u32 %rd15,%r0;
mul.lo.u64 %rd16,%rd7,%rd14;
setp.ge.u64 %p0,%rd15,%rd16;
@%p0 bra NB_DONE;
rem.u64 %rd17,%rd15,%rd14;
div.u64 %rd18,%rd15,%rd14;
div.u64 %rd19,%rd13,%rd14;
mul.lo.u64 %rd20,%rd11,%rd12;
mul.lo.u64 %rd21,%rd20,%rd19;
setp.ne.u32 %p1,%r1,0;
@%p1 bra NB_INIT_DONE;
st.shared.u32 [ns+0],0;
st.shared.u32 [ns+4],0;
NB_INIT_DONE:
bar.sync 0;
mul.lo.u64 %rd22,%rd15,8;
add.u64 %rd23,%rd6,%rd22;
ld.global.f32 %f0,[%rd23+0];
ld.global.f32 %f1,[%rd23+4];
cvt.u64.u32 %rd24,%r1;
mov.f32 %f18,0f;
mov.f32 %f19,0f;
NB_REDUCE:
setp.ge.u64 %p2,%rd24,%rd21;
@%p2 bra NB_REDUCE_DONE;
div.u64 %rd25,%rd24,%rd19;
rem.u64 %rd26,%rd24,%rd19;
mad.lo.u64 %rd27,%rd17,%rd19,%rd26;
mul.lo.u64 %rd28,%rd18,%rd20;
add.u64 %rd28,%rd28,%rd25;
mul.lo.u64 %rd28,%rd28,%rd13;
add.u64 %rd28,%rd28,%rd27;
shl.b64 %rd29,%rd28,1;
add.u64 %rd30,%rd0,%rd29;
ld.global.b16 %h0,[%rd30];
cvt.f32.bf16 %f2,%h0;
sub.f32 %f3,%f2,%f0;
mul.f32 %f3,%f3,%f1;
shl.b64 %rd31,%rd28,2;
add.u64 %rd32,%rd1,%rd31;
ld.global.f32 %f4,[%rd32];
shl.b64 %rd33,%rd27,1;
add.u64 %rd34,%rd3,%rd33;
ld.global.b16 %h1,[%rd34];
cvt.f32.bf16 %f5,%h1;
mul.f32 %f6,%f4,%f5;
mul.f32 %f7,%f6,%f3;
add.f32 %f18,%f18,%f6;
add.f32 %f19,%f19,%f7;
mul.f32 %f10,%f4,%f3;
shl.b64 %rd35,%rd27,2;
add.u64 %rd36,%rd4,%rd35;
add.u64 %rd37,%rd5,%rd35;
atom.global.add.f32 %f11,[%rd36],%f10;
atom.global.add.f32 %f12,[%rd37],%f4;
add.u64 %rd24,%rd24,%rd2;
bra NB_REDUCE;
NB_REDUCE_DONE:
atom.shared.add.f32 %f8,[ns+0],%f18;
atom.shared.add.f32 %f9,[ns+4],%f19;
bar.sync 0;
ld.shared.f32 %f13,[ns+0];
ld.shared.f32 %f14,[ns+4];
cvt.rn.f32.u64 %f15,%rd21;
cvt.u64.u32 %rd24,%r1;
NB_WRITE:
setp.ge.u64 %p3,%rd24,%rd21;
@%p3 bra NB_DONE;
div.u64 %rd25,%rd24,%rd19;
rem.u64 %rd26,%rd24,%rd19;
mad.lo.u64 %rd27,%rd17,%rd19,%rd26;
mul.lo.u64 %rd28,%rd18,%rd20;
add.u64 %rd28,%rd28,%rd25;
mul.lo.u64 %rd28,%rd28,%rd13;
add.u64 %rd28,%rd28,%rd27;
shl.b64 %rd29,%rd28,1;
add.u64 %rd30,%rd0,%rd29;
ld.global.b16 %h0,[%rd30];
cvt.f32.bf16 %f2,%h0;
sub.f32 %f3,%f2,%f0;
mul.f32 %f3,%f3,%f1;
shl.b64 %rd31,%rd28,2;
add.u64 %rd32,%rd1,%rd31;
ld.global.f32 %f4,[%rd32];
shl.b64 %rd33,%rd27,1;
add.u64 %rd34,%rd3,%rd33;
ld.global.b16 %h1,[%rd34];
cvt.f32.bf16 %f5,%h1;
mul.f32 %f6,%f4,%f5;
mul.f32 %f16,%f15,%f6;
sub.f32 %f16,%f16,%f13;
fma.rn.f32 %f16,-%f3,%f14,%f16;
mul.f32 %f16,%f16,%f1;
div.rn.f32 %f16,%f16,%f15;
add.u64 %rd38,%rd2,%rd31;
atom.global.add.f32 %f17,[%rd38],%f16;
add.u64 %rd24,%rd24,%rd2;
bra NB_WRITE;
NB_DONE:
ret;
}
.endgpu
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment