Preskočite na sadržaj

Prevođenje OpenCL C koda program prevoditeljem Clang

Clang je program prevoditelj za programske jezike slične C-u (C, C++, Objective C/C++, OpenCL C i CUDA C/C++) i dio je projekta LLVM.

Mi ćemo se u nastavku ograničiti na prevođenje jezika OpenCL C u LLVM-ovu srednju reprezentaciju (koja se zatim može prevesti u asemblerski kod grafičkog ili osnovnog procesora), ali proces prevođenja jezika CUDA C/C++ i HIP bio bi vrlo sličan.

Izvorni kod zrna

Uzmimo da imamo jednostavno zrno u OpenCL C-u koje sve elemente u polju in postavlja na dvostruku vrijednost i sprema u polje out. Kod je oblika:

__kernel void multiply_by_two(__global float *in, __global float *out)
{
    int index = get_global_id(0);
    out[index] = 2.0f * in[index];
}

Uočimo poziv funkcije get_global_id(); kad bismo ovaj kod zrna prevodili unutar programa, ta funkcija bi postojala u zaglavlju upravljačkog programa grafičkog procesora i bila bi uključena u proces prevođenja. Radi jednostavnosti, ovdje ćemo umjesto uključivanja zaglavlja sami deklarirati tu funkciju:

int get_global_id(int index);

__kernel void multiply_by_two(__global float *in, __global float *out)
{
    int index = get_global_id(0);
    out[index] = 2.0f * in[index];
}

Spremimo taj kod u datoteku kernel.cl.

Prvo se uvjerimo da imamo Clang instaliran i provjerimo njegovu verziju.

$ clang --version
Debian clang version 11.0.1-2
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /usr/bin

Prevođenje u asemblerski kod

U nastavku nas zanima proces prevođenja zrna, a ne njegovo izvođenje. Clang možemo pokrenuti korištenjem parametra -S tako da izvede samo prevođenje programskog koda u asemblerski kod, ali ne i povezivanje s bibliotekama i izradu izvršne datoteke. Bez dodatnih parametara kojima bismo naveli arhitekturu dobivamo asemblerski kod za arhitekturu x86-64 (AMD64 i Intel EM64T), koju vidimo navedenu iznad u dijelu Target.

Pokretanjem naredbe

$ clang -S kernel.cl

dobivamo datoteku kernel.s sadržaja:

        .text
        .file   "kernel.cl"
        .globl  multiply_by_two                 # -- Begin function multiply_by_two
        .p2align        4, 0x90
        .type   multiply_by_two,@function
multiply_by_two:                        # @multiply_by_two
        .cfi_startproc
# %bb.0:
        pushq   %rbp
        .cfi_def_cfa_offset 16
        .cfi_offset %rbp, -16
        movq    %rsp, %rbp
        .cfi_def_cfa_register %rbp
        pushq   %r14
        pushq   %rbx
        .cfi_offset %rbx, -32
        .cfi_offset %r14, -24
        movq    %rsi, %r14
        movq    %rdi, %rbx
        xorl    %edi, %edi
        callq   get_global_id
        cltq
        movss   (%rbx,%rax,4), %xmm0            # xmm0 = mem[0],zero,zero,zero
        addss   %xmm0, %xmm0
        movss   %xmm0, (%r14,%rax,4)
        popq    %rbx
        popq    %r14
        popq    %rbp
        .cfi_def_cfa %rsp, 8
        retq
.Lfunc_end0:
        .size   multiply_by_two, .Lfunc_end0-multiply_by_two
        .cfi_endproc
                                        # -- End function
        .ident  "Debian clang version 11.0.1-2"
        .section        ".note.GNU-stack","",@progbits
        .addrsig

Želimo li neku drugu arhitekturu, npr. MIPS, možemo je navesti parametrom -target:

$ clang -target mips -S kernel.cl

Datoteka kernel.s je sada sadržaja:

        .text
        .abicalls
        .option pic0
        .section        .mdebug.abi32,"",@progbits
        .nan    legacy
        .module fp=xx
        .module nooddspreg
        .text
        .file   "kernel.cl"
        .globl  multiply_by_two                 # -- Begin function multiply_by_two
        .p2align        2
        .type   multiply_by_two,@function
        .set    nomicromips
        .set    nomips16
        .ent    multiply_by_two
multiply_by_two:                        # @multiply_by_two
        .frame  $fp,32,$ra
        .mask   0xc0030000,-4
        .fmask  0x00000000,0
        .set    noreorder
        .set    nomacro
        .set    noat
# %bb.0:
        addiu   $sp, $sp, -32
        sw      $ra, 28($sp)                    # 4-byte Folded Spill
        sw      $fp, 24($sp)                    # 4-byte Folded Spill
        sw      $17, 20($sp)                    # 4-byte Folded Spill
        sw      $16, 16($sp)                    # 4-byte Folded Spill
        move    $fp, $sp
        move    $16, $5
        move    $17, $4
        jal     get_global_id
        addiu   $4, $zero, 0
        sll     $1, $2, 2
        lwxc1   $f0, $1($17)
        add.s   $f0, $f0, $f0
        swxc1   $f0, $1($16)
        move    $sp, $fp
        lw      $16, 16($sp)                    # 4-byte Folded Reload
        lw      $17, 20($sp)                    # 4-byte Folded Reload
        lw      $fp, 24($sp)                    # 4-byte Folded Reload
        lw      $ra, 28($sp)                    # 4-byte Folded Reload
        jr      $ra
        addiu   $sp, $sp, 32
        .set    at
        .set    macro
        .set    reorder
        .end    multiply_by_two
$func_end0:
        .size   multiply_by_two, ($func_end0)-multiply_by_two
                                        # -- End function
        .ident  "Debian clang version 11.0.1-2"
        .section        ".note.GNU-stack","",@progbits
        .addrsig
        .text

Prevođenjem u asemblerski kod za grafičke procesore bavit ćemo se malo kasnije pa ćemo i analizirati instrukcije koje se nalaze u kodu.

Prevođenje u srednju reprezentaciju

Clang možemo zaustaviti i prije nego proizvede asemblerski kod, odnosno reći mu da izvede samo prevođenje koda u LLVM-ovu srednju reprezentaciju dodavanjem parametra -emit-llvm.

$ clang -S -emit-llvm kernel.cl

Dobivamo datoteku kernel.ll sadržaja:

; ModuleID = 'kernel.cl'
source_filename = "kernel.cl"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-pc-linux-gnu"

; Function Attrs: convergent norecurse nounwind uwtable
define dso_local spir_kernel void @multiply_by_two(float* nocapture readonly %0, float* nocapture %1) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 {
%3 = tail call i32 @get_global_id(i32 0) #2
%4 = sext i32 %3 to i64
%5 = getelementptr inbounds float, float* %0, i64 %4
%6 = load float, float* %5, align 4, !tbaa !7
%7 = fmul float %6, 2.000000e+00
%8 = getelementptr inbounds float, float* %1, i64 %4
store float %7, float* %8, align 4, !tbaa !7
ret void
}

; Function Attrs: convergent
declare dso_local i32 @get_global_id(i32) local_unnamed_addr #1

attributes #0 = { convergent norecurse nounwind uwtable "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #2 = { convergent nounwind }

!llvm.module.flags = !{!0}
!opencl.ocl.version = !{!1}
!llvm.ident = !{!2}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 0}
!2 = !{!"Debian clang version 11.0.1-2"}
!3 = !{i32 1, i32 1}
!4 = !{!"none", !"none"}
!5 = !{!"float*", !"float*"}
!6 = !{!"", !""}
!7 = !{!8, !8, i64 0}
!8 = !{!"float", !9, i64 0}
!9 = !{!"omnipotent char", !10, i64 0}
!10 = !{!"Simple C/C++ TBAA"}

Uočimo da srednja reprezentacija sadrži atribute kao što su "target-cpu"="x86-64" "target-features"="+cx8,+fxsr,+mmx,+sse,+sse2,+x87" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" koji su specifični za x86-64 i bit će korišteni kod prevođenja u asemblerski kod. Srednju reprezentaciju za MIPS možemo dobiti naredbom clang -target mips -S -emit-llvm kernel.cl i ona je oblika:

; ModuleID = 'kernel.cl'
source_filename = "kernel.cl"
target datalayout = "E-m:m-p:32:32-i8:8:32-i16:16:32-i64:64-n32-S64"
target triple = "mips"

; Function Attrs: convergent norecurse nounwind
define dso_local spir_kernel void @multiply_by_two(float* nocapture readonly %0, float* nocapture %1) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 {
  %3 = tail call i32 @get_global_id(i32 signext 0) #2
  %4 = getelementptr inbounds float, float* %0, i32 %3
  %5 = load float, float* %4, align 4, !tbaa !7
  %6 = fmul float %5, 2.000000e+00
  %7 = getelementptr inbounds float, float* %1, i32 %3
  store float %6, float* %7, align 4, !tbaa !7
  ret void
}

; Function Attrs: convergent
declare dso_local i32 @get_global_id(i32 signext) local_unnamed_addr #1

attributes #0 = { convergent norecurse nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="mips32r2" "target-features"="+fpxx,+mips32r2,+nooddspreg,-noabicalls" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="mips32r2" "target-features"="+fpxx,+mips32r2,+nooddspreg,-noabicalls" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #2 = { convergent nounwind }

!llvm.module.flags = !{!0}
!opencl.ocl.version = !{!1}
!llvm.ident = !{!2}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 0}
!2 = !{!"Debian clang version 11.0.1-2"}
!3 = !{i32 1, i32 1}
!4 = !{!"none", !"none"}
!5 = !{!"float*", !"float*"}
!6 = !{!"", !""}
!7 = !{!8, !8, i64 0}
!8 = !{!"float", !9, i64 0}
!9 = !{!"omnipotent char", !10, i64 0}
!10 = !{!"Simple C/C++ TBAA"}

Srednjom reprezentacijom se nećemo dalje baviti, ali je načelno vrlo korisna za analizu načina rada optimizacija jer ima jednostavniju sintaksu od asemblerskih jezika. Ona se može prevesti u asemblerski kod LLVM-ovim alatom llc.

Author: Vedran Miletić