src/gallium/drivers/radeonsi/si_shader_llvm.c

   1 /*
   2  * Copyright 2016 Advanced Micro Devices, Inc.
   3  * All Rights Reserved.
   4  *
   5  * Permission is hereby granted, free of charge, to any person obtaining a
   6  * copy of this software and associated documentation files (the "Software"),
   7  * to deal in the Software without restriction, including without limitation
   8  * on the rights to use, copy, modify, merge, publish, distribute, sub
   9  * license, and/or sell copies of the Software, and to permit persons to whom
  10  * the Software is furnished to do so, subject to the following conditions:
  11  *
  12  * The above copyright notice and this permission notice (including the next
  13  * paragraph) shall be included in all copies or substantial portions of the
  14  * Software.
  15  *
  16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  18  * FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT. IN NO EVENT SHALL
  19  * THE AUTHOR(S) AND/OR THEIR SUPPLIERS BE LIABLE FOR ANY CLAIM,
  20  * DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR
  21  * OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE
  22  * USE OR OTHER DEALINGS IN THE SOFTWARE.
  23  */
  24
  25 #include "si_shader_internal.h"
  26 #include "si_pipe.h"
  27 #include "ac_rtld.h"
  28 #include "ac_nir_to_llvm.h"
  29 #include "sid.h"
  30
  31 #include "tgsi/tgsi_from_mesa.h"
  32 #include "util/u_memory.h"
  33
  34 struct si_llvm_diagnostics {
  35         struct pipe_debug_callback *debug;
  36         unsigned retval;
  37 };
  38
  39 static void si_diagnostic_handler(LLVMDiagnosticInfoRef di, void *context)
  40 {
  41         struct si_llvm_diagnostics *diag = (struct si_llvm_diagnostics *)context;
  42         LLVMDiagnosticSeverity severity = LLVMGetDiagInfoSeverity(di);
  43         const char *severity_str = NULL;
  44
  45         switch (severity) {
  46         case LLVMDSError:
  47                 severity_str = "error";
  48                 break;
  49         case LLVMDSWarning:
  50                 severity_str = "warning";
  51                 break;
  52         case LLVMDSRemark:
  53         case LLVMDSNote:
  54         default:
  55                 return;
  56         }
  57
  58         char *description = LLVMGetDiagInfoDescription(di);
  59
  60         pipe_debug_message(diag->debug, SHADER_INFO,
  61                            "LLVM diagnostic (%s): %s", severity_str, description);
  62
  63         if (severity == LLVMDSError) {
  64                 diag->retval = 1;
  65                 fprintf(stderr,"LLVM triggered Diagnostic Handler: %s\n", description);
  66         }
  67
  68         LLVMDisposeMessage(description);
  69 }
  70
  71 bool si_compile_llvm(struct si_screen *sscreen,
  72                      struct si_shader_binary *binary,
  73                      struct ac_shader_config *conf,
  74                      struct ac_llvm_compiler *compiler,
  75                      struct ac_llvm_context *ac,
  76                      struct pipe_debug_callback *debug,
  77                      enum pipe_shader_type shader_type,
  78                      const char *name,
  79                      bool less_optimized)
  80 {
  81         unsigned count = p_atomic_inc_return(&sscreen->num_compilations);
  82
  83         if (si_can_dump_shader(sscreen, shader_type)) {
  84                 fprintf(stderr, "radeonsi: Compiling shader %d\n", count);
  85
  86                 if (!(sscreen->debug_flags & (DBG(NO_IR) | DBG(PREOPT_IR)))) {
  87                         fprintf(stderr, "%s LLVM IR:\n\n", name);
  88                         ac_dump_module(ac->module);
  89                         fprintf(stderr, "\n");
  90                 }
  91         }
  92
  93         if (sscreen->record_llvm_ir) {
  94                 char *ir = LLVMPrintModuleToString(ac->module);
  95                 binary->llvm_ir_string = strdup(ir);
  96                 LLVMDisposeMessage(ir);
  97         }
  98
  99         if (!si_replace_shader(count, binary)) {
 100                 struct ac_compiler_passes *passes = compiler->passes;
 101
 102                 if (ac->wave_size == 32)
 103                         passes = compiler->passes_wave32;
 104                 else if (less_optimized && compiler->low_opt_passes)
 105                         passes = compiler->low_opt_passes;
 106
 107                 struct si_llvm_diagnostics diag = {debug};
 108                 LLVMContextSetDiagnosticHandler(ac->context, si_diagnostic_handler, &diag);
 109
 110                 if (!ac_compile_module_to_elf(passes, ac->module,
 111                                               (char **)&binary->elf_buffer,
 112                                               &binary->elf_size))
 113                         diag.retval = 1;
 114
 115                 if (diag.retval != 0) {
 116                         pipe_debug_message(debug, SHADER_INFO, "LLVM compilation failed");
 117                         return false;
 118                 }
 119         }
 120
 121         struct ac_rtld_binary rtld;
 122         if (!ac_rtld_open(&rtld, (struct ac_rtld_open_info){
 123                         .info = &sscreen->info,
 124                         .shader_type = tgsi_processor_to_shader_stage(shader_type),
 125                         .wave_size = ac->wave_size,
 126                         .num_parts = 1,
 127                         .elf_ptrs = &binary->elf_buffer,
 128                         .elf_sizes = &binary->elf_size }))
 129                 return false;
 130
 131         bool ok = ac_rtld_read_config(&rtld, conf);
 132         ac_rtld_close(&rtld);
 133         return ok;
 134 }
 135
 136 void si_llvm_context_init(struct si_shader_context *ctx,
 137                           struct si_screen *sscreen,
 138                           struct ac_llvm_compiler *compiler,
 139                           unsigned wave_size)
 140 {
 141         memset(ctx, 0, sizeof(*ctx));
 142         ctx->screen = sscreen;
 143         ctx->compiler = compiler;
 144
 145         ac_llvm_context_init(&ctx->ac, compiler, sscreen->info.chip_class,
 146                              sscreen->info.family,
 147                              AC_FLOAT_MODE_NO_SIGNED_ZEROS_FP_MATH,
 148                              wave_size, 64);
 149 }
 150
 151 void si_llvm_create_func(struct si_shader_context *ctx, const char *name,
 152                          LLVMTypeRef *return_types, unsigned num_return_elems,
 153                          unsigned max_workgroup_size)
 154 {
 155         LLVMTypeRef ret_type;
 156         enum ac_llvm_calling_convention call_conv;
 157         enum pipe_shader_type real_shader_type;
 158
 159         if (num_return_elems)
 160                 ret_type = LLVMStructTypeInContext(ctx->ac.context,
 161                                                    return_types,
 162                                                    num_return_elems, true);
 163         else
 164                 ret_type = ctx->ac.voidt;
 165
 166         real_shader_type = ctx->type;
 167
 168         /* LS is merged into HS (TCS), and ES is merged into GS. */
 169         if (ctx->screen->info.chip_class >= GFX9) {
 170                 if (ctx->shader->key.as_ls)
 171                         real_shader_type = PIPE_SHADER_TESS_CTRL;
 172                 else if (ctx->shader->key.as_es || ctx->shader->key.as_ngg)
 173                         real_shader_type = PIPE_SHADER_GEOMETRY;
 174         }
 175
 176         switch (real_shader_type) {
 177         case PIPE_SHADER_VERTEX:
 178         case PIPE_SHADER_TESS_EVAL:
 179                 call_conv = AC_LLVM_AMDGPU_VS;
 180                 break;
 181         case PIPE_SHADER_TESS_CTRL:
 182                 call_conv = AC_LLVM_AMDGPU_HS;
 183                 break;
 184         case PIPE_SHADER_GEOMETRY:
 185                 call_conv = AC_LLVM_AMDGPU_GS;
 186                 break;
 187         case PIPE_SHADER_FRAGMENT:
 188                 call_conv = AC_LLVM_AMDGPU_PS;
 189                 break;
 190         case PIPE_SHADER_COMPUTE:
 191                 call_conv = AC_LLVM_AMDGPU_CS;
 192                 break;
 193         default:
 194                 unreachable("Unhandle shader type");
 195         }
 196
 197         /* Setup the function */
 198         ctx->return_type = ret_type;
 199         ctx->main_fn = ac_build_main(&ctx->args, &ctx->ac, call_conv, name,
 200                                      ret_type, ctx->ac.module);
 201         ctx->return_value = LLVMGetUndef(ctx->return_type);
 202
 203         if (ctx->screen->info.address32_hi) {
 204                 ac_llvm_add_target_dep_function_attr(ctx->main_fn,
 205                                                      "amdgpu-32bit-address-high-bits",
 206                                                      ctx->screen->info.address32_hi);
 207         }
 208
 209         LLVMAddTargetDependentFunctionAttr(ctx->main_fn,
 210                                            "no-signed-zeros-fp-math",
 211                                            "true");
 212
 213         ac_llvm_set_workgroup_size(ctx->main_fn, max_workgroup_size);
 214 }
 215
 216 void si_llvm_optimize_module(struct si_shader_context *ctx)
 217 {
 218         /* Dump LLVM IR before any optimization passes */
 219         if (ctx->screen->debug_flags & DBG(PREOPT_IR) &&
 220             si_can_dump_shader(ctx->screen, ctx->type))
 221                 LLVMDumpModule(ctx->ac.module);
 222
 223         /* Run the pass */
 224         LLVMRunPassManager(ctx->compiler->passmgr, ctx->ac.module);
 225         LLVMDisposeBuilder(ctx->ac.builder);
 226 }
 227
 228 void si_llvm_dispose(struct si_shader_context *ctx)
 229 {
 230         LLVMDisposeModule(ctx->ac.module);
 231         LLVMContextDispose(ctx->ac.context);
 232         ac_llvm_context_dispose(&ctx->ac);
 233 }
 234
 235 /**
 236  * Load a dword from a constant buffer.
 237  */
 238 LLVMValueRef si_buffer_load_const(struct si_shader_context *ctx,
 239                                   LLVMValueRef resource, LLVMValueRef offset)
 240 {
 241         return ac_build_buffer_load(&ctx->ac, resource, 1, NULL, offset, NULL,
 242                                     0, 0, true, true);
 243 }
 244
 245 void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret)
 246 {
 247         if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
 248                 LLVMBuildRetVoid(ctx->ac.builder);
 249         else
 250                 LLVMBuildRet(ctx->ac.builder, ret);
 251 }
 252
 253 LLVMValueRef si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret,
 254                                  struct ac_arg param, unsigned return_index)
 255 {
 256         return LLVMBuildInsertValue(ctx->ac.builder, ret,
 257                                     ac_get_arg(&ctx->ac, param),
 258                                     return_index, "");
 259 }
 260
 261 LLVMValueRef si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret,
 262                                        struct ac_arg param, unsigned return_index)
 263 {
 264         LLVMBuilderRef builder = ctx->ac.builder;
 265         LLVMValueRef p = ac_get_arg(&ctx->ac, param);
 266
 267         return LLVMBuildInsertValue(builder, ret,
 268                                     ac_to_float(&ctx->ac, p),
 269                                     return_index, "");
 270 }
 271
 272 LLVMValueRef si_insert_input_ptr(struct si_shader_context *ctx, LLVMValueRef ret,
 273                                  struct ac_arg param, unsigned return_index)
 274 {
 275         LLVMBuilderRef builder = ctx->ac.builder;
 276         LLVMValueRef ptr = ac_get_arg(&ctx->ac, param);
 277         ptr = LLVMBuildPtrToInt(builder, ptr, ctx->ac.i32, "");
 278         return LLVMBuildInsertValue(builder, ret, ptr, return_index, "");
 279 }
 280
 281 LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
 282 {
 283         LLVMValueRef ptr[2], list;
 284         bool merged_shader = si_is_merged_shader(ctx->shader);
 285
 286         ptr[0] = LLVMGetParam(ctx->main_fn, (merged_shader ? 8 : 0) + SI_SGPR_RW_BUFFERS);
 287         list = LLVMBuildIntToPtr(ctx->ac.builder, ptr[0],
 288                                  ac_array_in_const32_addr_space(ctx->ac.v4i32), "");
 289         return list;
 290 }
 291
 292 LLVMValueRef si_build_gather_64bit(struct si_shader_context *ctx,
 293                                    LLVMTypeRef type, LLVMValueRef val1,
 294                                    LLVMValueRef val2)
 295 {
 296         LLVMValueRef values[2] = {
 297                 ac_to_integer(&ctx->ac, val1),
 298                 ac_to_integer(&ctx->ac, val2),
 299         };
 300         LLVMValueRef result = ac_build_gather_values(&ctx->ac, values, 2);
 301         return LLVMBuildBitCast(ctx->ac.builder, result, type, "");
 302 }
 303
 304 void si_llvm_emit_barrier(struct si_shader_context *ctx)
 305 {
 306         /* GFX6 only (thanks to a hw bug workaround):
 307          * The real barrier instruction isn’t needed, because an entire patch
 308          * always fits into a single wave.
 309          */
 310         if (ctx->screen->info.chip_class == GFX6 &&
 311             ctx->type == PIPE_SHADER_TESS_CTRL) {
 312                 ac_build_waitcnt(&ctx->ac, AC_WAIT_LGKM | AC_WAIT_VLOAD | AC_WAIT_VSTORE);
 313                 return;
 314         }
 315
 316         ac_build_s_barrier(&ctx->ac);
 317 }
 318
 319 /* Ensure that the esgs ring is declared.
 320  *
 321  * We declare it with 64KB alignment as a hint that the
 322  * pointer value will always be 0.
 323  */
 324 void si_llvm_declare_esgs_ring(struct si_shader_context *ctx)
 325 {
 326         if (ctx->esgs_ring)
 327                 return;
 328
 329         assert(!LLVMGetNamedGlobal(ctx->ac.module, "esgs_ring"));
 330
 331         ctx->esgs_ring = LLVMAddGlobalInAddressSpace(
 332                 ctx->ac.module, LLVMArrayType(ctx->ac.i32, 0),
 333                 "esgs_ring",
 334                 AC_ADDR_SPACE_LDS);
 335         LLVMSetLinkage(ctx->esgs_ring, LLVMExternalLinkage);
 336         LLVMSetAlignment(ctx->esgs_ring, 64 * 1024);
 337 }
 338
 339 void si_init_exec_from_input(struct si_shader_context *ctx, struct ac_arg param,
 340                              unsigned bitoffset)
 341 {
 342         LLVMValueRef args[] = {
 343                 ac_get_arg(&ctx->ac, param),
 344                 LLVMConstInt(ctx->ac.i32, bitoffset, 0),
 345         };
 346         ac_build_intrinsic(&ctx->ac,
 347                            "llvm.amdgcn.init.exec.from.input",
 348                            ctx->ac.voidt, args, 2, AC_FUNC_ATTR_CONVERGENT);
 349 }
 350
 351 /**
 352  * Get the value of a shader input parameter and extract a bitfield.
 353  */
 354 static LLVMValueRef unpack_llvm_param(struct si_shader_context *ctx,
 355                                       LLVMValueRef value, unsigned rshift,
 356                                       unsigned bitwidth)
 357 {
 358         if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
 359                 value = ac_to_integer(&ctx->ac, value);
 360
 361         if (rshift)
 362                 value = LLVMBuildLShr(ctx->ac.builder, value,
 363                                       LLVMConstInt(ctx->ac.i32, rshift, 0), "");
 364
 365         if (rshift + bitwidth < 32) {
 366                 unsigned mask = (1 << bitwidth) - 1;
 367                 value = LLVMBuildAnd(ctx->ac.builder, value,
 368                                      LLVMConstInt(ctx->ac.i32, mask, 0), "");
 369         }
 370
 371         return value;
 372 }
 373
 374 LLVMValueRef si_unpack_param(struct si_shader_context *ctx,
 375                              struct ac_arg param, unsigned rshift,
 376                              unsigned bitwidth)
 377 {
 378         LLVMValueRef value = ac_get_arg(&ctx->ac, param);
 379
 380         return unpack_llvm_param(ctx, value, rshift, bitwidth);
 381 }
 382
 383 LLVMValueRef si_get_primitive_id(struct si_shader_context *ctx,
 384                                  unsigned swizzle)
 385 {
 386         if (swizzle > 0)
 387                 return ctx->ac.i32_0;
 388
 389         switch (ctx->type) {
 390         case PIPE_SHADER_VERTEX:
 391                 return ac_get_arg(&ctx->ac, ctx->vs_prim_id);
 392         case PIPE_SHADER_TESS_CTRL:
 393                 return ac_get_arg(&ctx->ac, ctx->args.tcs_patch_id);
 394         case PIPE_SHADER_TESS_EVAL:
 395                 return ac_get_arg(&ctx->ac, ctx->args.tes_patch_id);
 396         case PIPE_SHADER_GEOMETRY:
 397                 return ac_get_arg(&ctx->ac, ctx->args.gs_prim_id);
 398         default:
 399                 assert(0);
 400                 return ctx->ac.i32_0;
 401         }
 402 }
 403
 404 LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
 405 {
 406         struct si_shader_context *ctx = si_shader_context_from_abi(abi);
 407
 408         LLVMValueRef values[3];
 409         LLVMValueRef result;
 410         unsigned i;
 411         unsigned *properties = ctx->shader->selector->info.properties;
 412
 413         if (properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] != 0) {
 414                 unsigned sizes[3] = {
 415                         properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH],
 416                         properties[TGSI_PROPERTY_CS_FIXED_BLOCK_HEIGHT],
 417                         properties[TGSI_PROPERTY_CS_FIXED_BLOCK_DEPTH]
 418                 };
 419
 420                 for (i = 0; i < 3; ++i)
 421                         values[i] = LLVMConstInt(ctx->ac.i32, sizes[i], 0);
 422
 423                 result = ac_build_gather_values(&ctx->ac, values, 3);
 424         } else {
 425                 result = ac_get_arg(&ctx->ac, ctx->block_size);
 426         }
 427
 428         return result;
 429 }
 430
 431 void si_llvm_declare_compute_memory(struct si_shader_context *ctx)
 432 {
 433         struct si_shader_selector *sel = ctx->shader->selector;
 434         unsigned lds_size = sel->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE];
 435
 436         LLVMTypeRef i8p = LLVMPointerType(ctx->ac.i8, AC_ADDR_SPACE_LDS);
 437         LLVMValueRef var;
 438
 439         assert(!ctx->ac.lds);
 440
 441         var = LLVMAddGlobalInAddressSpace(ctx->ac.module,
 442                                           LLVMArrayType(ctx->ac.i8, lds_size),
 443                                           "compute_lds",
 444                                           AC_ADDR_SPACE_LDS);
 445         LLVMSetAlignment(var, 64 * 1024);
 446
 447         ctx->ac.lds = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
 448 }
 449
 450 bool si_nir_build_llvm(struct si_shader_context *ctx, struct nir_shader *nir)
 451 {
 452         if (nir->info.stage == MESA_SHADER_VERTEX) {
 453                 si_llvm_load_vs_inputs(ctx, nir);
 454         } else if (nir->info.stage == MESA_SHADER_FRAGMENT) {
 455                 unsigned colors_read =
 456                         ctx->shader->selector->info.colors_read;
 457                 LLVMValueRef main_fn = ctx->main_fn;
 458
 459                 LLVMValueRef undef = LLVMGetUndef(ctx->ac.f32);
 460
 461                 unsigned offset = SI_PARAM_POS_FIXED_PT + 1;
 462
 463                 if (colors_read & 0x0f) {
 464                         unsigned mask = colors_read & 0x0f;
 465                         LLVMValueRef values[4];
 466                         values[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : undef;
 467                         values[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : undef;
 468                         values[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : undef;
 469                         values[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : undef;
 470                         ctx->abi.color0 =
 471                                 ac_to_integer(&ctx->ac,
 472                                               ac_build_gather_values(&ctx->ac, values, 4));
 473                 }
 474                 if (colors_read & 0xf0) {
 475                         unsigned mask = (colors_read & 0xf0) >> 4;
 476                         LLVMValueRef values[4];
 477                         values[0] = mask & 0x1 ? LLVMGetParam(main_fn, offset++) : undef;
 478                         values[1] = mask & 0x2 ? LLVMGetParam(main_fn, offset++) : undef;
 479                         values[2] = mask & 0x4 ? LLVMGetParam(main_fn, offset++) : undef;
 480                         values[3] = mask & 0x8 ? LLVMGetParam(main_fn, offset++) : undef;
 481                         ctx->abi.color1 =
 482                                 ac_to_integer(&ctx->ac,
 483                                               ac_build_gather_values(&ctx->ac, values, 4));
 484                 }
 485
 486                 ctx->abi.interp_at_sample_force_center =
 487                         ctx->shader->key.mono.u.ps.interpolate_at_sample_force_center;
 488         } else if (nir->info.stage == MESA_SHADER_COMPUTE) {
 489                 if (nir->info.cs.user_data_components_amd) {
 490                         ctx->abi.user_data = ac_get_arg(&ctx->ac, ctx->cs_user_data);
 491                         ctx->abi.user_data = ac_build_expand_to_vec4(&ctx->ac, ctx->abi.user_data,
 492                                                                      nir->info.cs.user_data_components_amd);
 493                 }
 494         }
 495
 496         ctx->abi.inputs = &ctx->inputs[0];
 497         ctx->abi.clamp_shadow_reference = true;
 498         ctx->abi.robust_buffer_access = true;
 499
 500         if (ctx->shader->selector->info.properties[TGSI_PROPERTY_CS_LOCAL_SIZE]) {
 501                 assert(gl_shader_stage_is_compute(nir->info.stage));
 502                 si_llvm_declare_compute_memory(ctx);
 503         }
 504         ac_nir_translate(&ctx->ac, &ctx->abi, &ctx->args, nir);
 505
 506         return true;
 507 }
 508
 509 /**
 510  * Given a list of shader part functions, build a wrapper function that
 511  * runs them in sequence to form a monolithic shader.
 512  */
 513 void si_build_wrapper_function(struct si_shader_context *ctx, LLVMValueRef *parts,
 514                                unsigned num_parts, unsigned main_part,
 515                                unsigned next_shader_first_part)
 516 {
 517         LLVMBuilderRef builder = ctx->ac.builder;
 518         /* PS epilog has one arg per color component; gfx9 merged shader
 519          * prologs need to forward 40 SGPRs.
 520          */
 521         LLVMValueRef initial[AC_MAX_ARGS], out[AC_MAX_ARGS];
 522         LLVMTypeRef function_type;
 523         unsigned num_first_params;
 524         unsigned num_out, initial_num_out;
 525         ASSERTED unsigned num_out_sgpr; /* used in debug checks */
 526         ASSERTED unsigned initial_num_out_sgpr; /* used in debug checks */
 527         unsigned num_sgprs, num_vgprs;
 528         unsigned gprs;
 529
 530         memset(&ctx->args, 0, sizeof(ctx->args));
 531
 532         for (unsigned i = 0; i < num_parts; ++i) {
 533                 ac_add_function_attr(ctx->ac.context, parts[i], -1,
 534                                      AC_FUNC_ATTR_ALWAYSINLINE);
 535                 LLVMSetLinkage(parts[i], LLVMPrivateLinkage);
 536         }
 537
 538         /* The parameters of the wrapper function correspond to those of the
 539          * first part in terms of SGPRs and VGPRs, but we use the types of the
 540          * main part to get the right types. This is relevant for the
 541          * dereferenceable attribute on descriptor table pointers.
 542          */
 543         num_sgprs = 0;
 544         num_vgprs = 0;
 545
 546         function_type = LLVMGetElementType(LLVMTypeOf(parts[0]));
 547         num_first_params = LLVMCountParamTypes(function_type);
 548
 549         for (unsigned i = 0; i < num_first_params; ++i) {
 550                 LLVMValueRef param = LLVMGetParam(parts[0], i);
 551
 552                 if (ac_is_sgpr_param(param)) {
 553                         assert(num_vgprs == 0);
 554                         num_sgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
 555                 } else {
 556                         num_vgprs += ac_get_type_size(LLVMTypeOf(param)) / 4;
 557                 }
 558         }
 559
 560         gprs = 0;
 561         while (gprs < num_sgprs + num_vgprs) {
 562                 LLVMValueRef param = LLVMGetParam(parts[main_part], ctx->args.arg_count);
 563                 LLVMTypeRef type = LLVMTypeOf(param);
 564                 unsigned size = ac_get_type_size(type) / 4;
 565
 566                 /* This is going to get casted anyways, so we don't have to
 567                  * have the exact same type. But we do have to preserve the
 568                  * pointer-ness so that LLVM knows about it.
 569                  */
 570                 enum ac_arg_type arg_type = AC_ARG_INT;
 571                 if (LLVMGetTypeKind(type) == LLVMPointerTypeKind) {
 572                         type = LLVMGetElementType(type);
 573
 574                         if (LLVMGetTypeKind(type) == LLVMVectorTypeKind) {
 575                                 if (LLVMGetVectorSize(type) == 4)
 576                                         arg_type = AC_ARG_CONST_DESC_PTR;
 577                                 else if (LLVMGetVectorSize(type) == 8)
 578                                         arg_type = AC_ARG_CONST_IMAGE_PTR;
 579                                 else
 580                                         assert(0);
 581                         } else if (type == ctx->ac.f32) {
 582                                 arg_type = AC_ARG_CONST_FLOAT_PTR;
 583                         } else {
 584                                 assert(0);
 585                         }
 586                 }
 587
 588                 ac_add_arg(&ctx->args, gprs < num_sgprs ? AC_ARG_SGPR : AC_ARG_VGPR,
 589                            size, arg_type, NULL);
 590
 591                 assert(ac_is_sgpr_param(param) == (gprs < num_sgprs));
 592                 assert(gprs + size <= num_sgprs + num_vgprs &&
 593                        (gprs >= num_sgprs || gprs + size <= num_sgprs));
 594
 595                 gprs += size;
 596         }
 597
 598         /* Prepare the return type. */
 599         unsigned num_returns = 0;
 600         LLVMTypeRef returns[AC_MAX_ARGS], last_func_type, return_type;
 601
 602         last_func_type = LLVMGetElementType(LLVMTypeOf(parts[num_parts - 1]));
 603         return_type = LLVMGetReturnType(last_func_type);
 604
 605         switch (LLVMGetTypeKind(return_type)) {
 606         case LLVMStructTypeKind:
 607                 num_returns = LLVMCountStructElementTypes(return_type);
 608                 assert(num_returns <= ARRAY_SIZE(returns));
 609                 LLVMGetStructElementTypes(return_type, returns);
 610                 break;
 611         case LLVMVoidTypeKind:
 612                 break;
 613         default:
 614                 unreachable("unexpected type");
 615         }
 616
 617         si_llvm_create_func(ctx, "wrapper", returns, num_returns,
 618                             si_get_max_workgroup_size(ctx->shader));
 619
 620         if (si_is_merged_shader(ctx->shader))
 621                 ac_init_exec_full_mask(&ctx->ac);
 622
 623         /* Record the arguments of the function as if they were an output of
 624          * a previous part.
 625          */
 626         num_out = 0;
 627         num_out_sgpr = 0;
 628
 629         for (unsigned i = 0; i < ctx->args.arg_count; ++i) {
 630                 LLVMValueRef param = LLVMGetParam(ctx->main_fn, i);
 631                 LLVMTypeRef param_type = LLVMTypeOf(param);
 632                 LLVMTypeRef out_type = ctx->args.args[i].file == AC_ARG_SGPR ? ctx->ac.i32 : ctx->ac.f32;
 633                 unsigned size = ac_get_type_size(param_type) / 4;
 634
 635                 if (size == 1) {
 636                         if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
 637                                 param = LLVMBuildPtrToInt(builder, param, ctx->ac.i32, "");
 638                                 param_type = ctx->ac.i32;
 639                         }
 640
 641                         if (param_type != out_type)
 642                                 param = LLVMBuildBitCast(builder, param, out_type, "");
 643                         out[num_out++] = param;
 644                 } else {
 645                         LLVMTypeRef vector_type = LLVMVectorType(out_type, size);
 646
 647                         if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
 648                                 param = LLVMBuildPtrToInt(builder, param, ctx->ac.i64, "");
 649                                 param_type = ctx->ac.i64;
 650                         }
 651
 652                         if (param_type != vector_type)
 653                                 param = LLVMBuildBitCast(builder, param, vector_type, "");
 654
 655                         for (unsigned j = 0; j < size; ++j)
 656                                 out[num_out++] = LLVMBuildExtractElement(
 657                                         builder, param, LLVMConstInt(ctx->ac.i32, j, 0), "");
 658                 }
 659
 660                 if (ctx->args.args[i].file == AC_ARG_SGPR)
 661                         num_out_sgpr = num_out;
 662         }
 663
 664         memcpy(initial, out, sizeof(out));
 665         initial_num_out = num_out;
 666         initial_num_out_sgpr = num_out_sgpr;
 667
 668         /* Now chain the parts. */
 669         LLVMValueRef ret = NULL;
 670         for (unsigned part = 0; part < num_parts; ++part) {
 671                 LLVMValueRef in[AC_MAX_ARGS];
 672                 LLVMTypeRef ret_type;
 673                 unsigned out_idx = 0;
 674                 unsigned num_params = LLVMCountParams(parts[part]);
 675
 676                 /* Merged shaders are executed conditionally depending
 677                  * on the number of enabled threads passed in the input SGPRs. */
 678                 if (si_is_multi_part_shader(ctx->shader) && part == 0) {
 679                         LLVMValueRef ena, count = initial[3];
 680
 681                         count = LLVMBuildAnd(builder, count,
 682                                              LLVMConstInt(ctx->ac.i32, 0x7f, 0), "");
 683                         ena = LLVMBuildICmp(builder, LLVMIntULT,
 684                                             ac_get_thread_id(&ctx->ac), count, "");
 685                         ac_build_ifcc(&ctx->ac, ena, 6506);
 686                 }
 687
 688                 /* Derive arguments for the next part from outputs of the
 689                  * previous one.
 690                  */
 691                 for (unsigned param_idx = 0; param_idx < num_params; ++param_idx) {
 692                         LLVMValueRef param;
 693                         LLVMTypeRef param_type;
 694                         bool is_sgpr;
 695                         unsigned param_size;
 696                         LLVMValueRef arg = NULL;
 697
 698                         param = LLVMGetParam(parts[part], param_idx);
 699                         param_type = LLVMTypeOf(param);
 700                         param_size = ac_get_type_size(param_type) / 4;
 701                         is_sgpr = ac_is_sgpr_param(param);
 702
 703                         if (is_sgpr) {
 704                                 ac_add_function_attr(ctx->ac.context, parts[part],
 705                                                      param_idx + 1, AC_FUNC_ATTR_INREG);
 706                         } else if (out_idx < num_out_sgpr) {
 707                                 /* Skip returned SGPRs the current part doesn't
 708                                  * declare on the input. */
 709                                 out_idx = num_out_sgpr;
 710                         }
 711
 712                         assert(out_idx + param_size <= (is_sgpr ? num_out_sgpr : num_out));
 713
 714                         if (param_size == 1)
 715                                 arg = out[out_idx];
 716                         else
 717                                 arg = ac_build_gather_values(&ctx->ac, &out[out_idx], param_size);
 718
 719                         if (LLVMTypeOf(arg) != param_type) {
 720                                 if (LLVMGetTypeKind(param_type) == LLVMPointerTypeKind) {
 721                                         if (LLVMGetPointerAddressSpace(param_type) ==
 722                                             AC_ADDR_SPACE_CONST_32BIT) {
 723                                                 arg = LLVMBuildBitCast(builder, arg, ctx->ac.i32, "");
 724                                                 arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
 725                                         } else {
 726                                                 arg = LLVMBuildBitCast(builder, arg, ctx->ac.i64, "");
 727                                                 arg = LLVMBuildIntToPtr(builder, arg, param_type, "");
 728                                         }
 729                                 } else {
 730                                         arg = LLVMBuildBitCast(builder, arg, param_type, "");
 731                                 }
 732                         }
 733
 734                         in[param_idx] = arg;
 735                         out_idx += param_size;
 736                 }
 737
 738                 ret = ac_build_call(&ctx->ac, parts[part], in, num_params);
 739
 740                 if (si_is_multi_part_shader(ctx->shader) &&
 741                     part + 1 == next_shader_first_part) {
 742                         ac_build_endif(&ctx->ac, 6506);
 743
 744                         /* The second half of the merged shader should use
 745                          * the inputs from the toplevel (wrapper) function,
 746                          * not the return value from the last call.
 747                          *
 748                          * That's because the last call was executed condi-
 749                          * tionally, so we can't consume it in the main
 750                          * block.
 751                          */
 752                         memcpy(out, initial, sizeof(initial));
 753                         num_out = initial_num_out;
 754                         num_out_sgpr = initial_num_out_sgpr;
 755                         continue;
 756                 }
 757
 758                 /* Extract the returned GPRs. */
 759                 ret_type = LLVMTypeOf(ret);
 760                 num_out = 0;
 761                 num_out_sgpr = 0;
 762
 763                 if (LLVMGetTypeKind(ret_type) != LLVMVoidTypeKind) {
 764                         assert(LLVMGetTypeKind(ret_type) == LLVMStructTypeKind);
 765
 766                         unsigned ret_size = LLVMCountStructElementTypes(ret_type);
 767
 768                         for (unsigned i = 0; i < ret_size; ++i) {
 769                                 LLVMValueRef val =
 770                                         LLVMBuildExtractValue(builder, ret, i, "");
 771
 772                                 assert(num_out < ARRAY_SIZE(out));
 773                                 out[num_out++] = val;
 774
 775                                 if (LLVMTypeOf(val) == ctx->ac.i32) {
 776                                         assert(num_out_sgpr + 1 == num_out);
 777                                         num_out_sgpr = num_out;
 778                                 }
 779                         }
 780                 }
 781         }
 782
 783         /* Return the value from the last part. */
 784         if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
 785                 LLVMBuildRetVoid(builder);
 786         else
 787                 LLVMBuildRet(builder, ret);
 788 }