xbc: An extensible MLIR compiler infrastructure for programming language and heterogeneous computing research
Creators: name: Anonymous
Publisher: name: Zenodo
Year: 2024

description: Abstract

This paper presents the design and development of the eXtensible base compiler (xbc) infrastructure, an extensible MLIR infrastructure for programming language and heterogeneous computing research. This extensible infrastructure permits testing novel programming languages and programming model concepts, such as support for quantum targets or alternative parallel programming models. It also provides a platform for compiler practitioners to meet the demands of a rapidly changing hardware landscape. xbc targets traditional X86 processors, GPUs (NVIDIA and AMD), and quantum-classical heterogeneous systems. Our results demonstrate comparable performance to vendor compilers for GPUs. 

Highlighting one of our results, we observe that NAS CG xbc's GPU version on NVIDIA A100 is 2.26x faster than Clang's OpenMP offload and 1.3x faster than NVIDIA's OpenACC compilers. For AMD MI250x, xbc is 5x faster than Clang's OpenMP offload and 5.5x faster than AMD Clang's OpenMP offload.

Files

The JSON files include the raw execution times for the plots in the paper.

Change made in CG on Frontier

* Original loop

      #pragma omp target map(rowstr[:0],colidx[:0],a[:0],p[:0],q[:0])
      #pragma omp teams distribute parallel for private(tmp1,tmp2,sum)
      for (j = 0; j < end; j++) {
         tmp1 = rowstr[j];
         tmp2 = rowstr[j+1];
         sum = 0.0;
         #pragma omp simd reduction(+:sum) private(tmp3)
         for (k = tmp1; k < tmp2; k++) {
            tmp3 = colidx[k];
            sum = sum + a[k]*p[tmp3];
         }
         q[j] = sum;
      }

* Modified loop

      #pragma omp target map(rowstr[:0],colidx[:0],a[:0],p[:0],q[:0])
      #pragma omp teams distribute private(tmp1,tmp2,sum)
      for (j = 0; j < end; j++) {
         tmp1 = rowstr[j];
         tmp2 = rowstr[j+1];
         sum = 0.0;
         #pragma omp parallel for simd reduction(+:sum) private(tmp3)
         for (k = tmp1; k < tmp2; k++) {
            tmp3 = colidx[k];
            sum = sum + a[k]*p[tmp3];
         }
         q[j] = sum;
      }

xb ops in xbc

// Control-flow operations
xb.at_exit xb.break xb.continue xb.generic_control_flow
xb.if xb.loop xb.range_for xb.return
xb.scope xb.yield
// Type-related operations
xb.addr_of xb.from_memref xb.get_member xb.global
xb.load xb.ptradd xb.ptr_cast xb.sizeof
xb.store xb.to_memref xb.var
// Generalizations of upstream MLIR
xb.arith_cast xb.arith_op xb.call xb.cast
xb.constant xb.func xb.math xb.select
xb.variant

par ops in xbc

par.async par.atomic par.barrier par.data_region
par.dealloc_stack par.dim par.get_stack_entry par.id
par.isolated_region par.map par.stack_mem par.reduce
par.region par.region_value par.wait par.yield

Concepts in xlg

// Symbol related concepts
xlg.symbol xlg.sym_table
// Concepts with attributes or types
xlg.attr xlg.typed
// Type concepts
xlg.type xlg.builtin_type xlg.derived_type
// Declaration concepts
xlg.decl xlg.named_decl xlg.value_decl xlg.type_decl
xlg.var_decl xlg.func_decl xlg.object_decl
// Statement concepts
xlg.stmt xlg.control_flow xlg.loop xlg.cf_terminator
// Expression concepts
xlg.expr xlg.const_expr xlg.cast_expr xlg.ref_expr

Constructs in xlg

// Types
xlg.array_type xlg.builtin_type xlg.ptr_type xlg.ref_type
xlg.rm_ref_type xlg.template_type xlg.type_of xlg.type_ref
// Declarations
xlg.func_context xlg.func_decl xlg.object_decl xlg.var_decl
// Statements
xlg.compound_stmt xlg.if_stmt xlg.range_for_stmt xlg.return_stmt
xlg.while_stmt xlg.yield_stmt
// Expressions
xlg.array_expr xlg.bin_op_expr xlg.call_expr xlg.cast_expr
xlg.const_expr xlg.from_xlg_expr xlg.list_expr xlg.load_expr
xlg.member_ref_expr xlg.range_expr xlg.ref_expr xlg.select_expr
xlg.store_expr xlg.sym_expr xlg.template_expr xlg.to_xlg_expr
xlg.unary_expr xlg.value_ref_expr