[Mesa-dev] [PATCH 05/11] radeonsi: use ctx->ac.builder
Marek Olšák
maraeo at gmail.com
Fri Sep 29 14:49:49 UTC 2017
From: Marek Olšák <marek.olsak at amd.com>
---
src/gallium/drivers/radeonsi/si_shader.c | 401 ++++++++++-----------
src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c | 161 ++++-----
src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c | 114 +++---
.../drivers/radeonsi/si_shader_tgsi_setup.c | 84 ++---
4 files changed, 341 insertions(+), 419 deletions(-)
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 88c3e56..74c11d3 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -247,34 +247,33 @@ build_phi(struct ac_llvm_context *ctx, LLVMTypeRef type,
return phi;
}
/**
* Get the value of a shader input parameter and extract a bitfield.
*/
static LLVMValueRef unpack_param(struct si_shader_context *ctx,
unsigned param, unsigned rshift,
unsigned bitwidth)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value = LLVMGetParam(ctx->main_fn,
param);
if (LLVMGetTypeKind(LLVMTypeOf(value)) == LLVMFloatTypeKind)
value = ac_to_integer(&ctx->ac, value);
if (rshift)
- value = LLVMBuildLShr(gallivm->builder, value,
+ value = LLVMBuildLShr(ctx->ac.builder, value,
LLVMConstInt(ctx->i32, rshift, 0), "");
if (rshift + bitwidth < 32) {
unsigned mask = (1 << bitwidth) - 1;
- value = LLVMBuildAnd(gallivm->builder, value,
+ value = LLVMBuildAnd(ctx->ac.builder, value,
LLVMConstInt(ctx->i32, mask, 0), "");
}
return value;
}
static LLVMValueRef get_rel_patch_id(struct si_shader_context *ctx)
{
switch (ctx->type) {
case PIPE_SHADER_TESS_CTRL:
@@ -364,52 +363,49 @@ get_tcs_out_patch0_patch_data_offset(struct si_shader_context *ctx)
return lp_build_mul_imm(&ctx->bld_base.uint_bld,
unpack_param(ctx,
ctx->param_tcs_out_lds_offsets,
16, 16),
4);
}
static LLVMValueRef
get_tcs_in_current_patch_offset(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef patch_stride = get_tcs_in_patch_stride(ctx);
LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
- return LLVMBuildMul(gallivm->builder, patch_stride, rel_patch_id, "");
+ return LLVMBuildMul(ctx->ac.builder, patch_stride, rel_patch_id, "");
}
static LLVMValueRef
get_tcs_out_current_patch_offset(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef patch0_offset = get_tcs_out_patch0_offset(ctx);
LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
- return LLVMBuildAdd(gallivm->builder, patch0_offset,
- LLVMBuildMul(gallivm->builder, patch_stride,
+ return LLVMBuildAdd(ctx->ac.builder, patch0_offset,
+ LLVMBuildMul(ctx->ac.builder, patch_stride,
rel_patch_id, ""),
"");
}
static LLVMValueRef
get_tcs_out_current_patch_data_offset(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef patch0_patch_data_offset =
get_tcs_out_patch0_patch_data_offset(ctx);
LLVMValueRef patch_stride = get_tcs_out_patch_stride(ctx);
LLVMValueRef rel_patch_id = get_rel_patch_id(ctx);
- return LLVMBuildAdd(gallivm->builder, patch0_patch_data_offset,
- LLVMBuildMul(gallivm->builder, patch_stride,
+ return LLVMBuildAdd(ctx->ac.builder, patch0_patch_data_offset,
+ LLVMBuildMul(ctx->ac.builder, patch_stride,
rel_patch_id, ""),
"");
}
static LLVMValueRef get_num_tcs_out_vertices(struct si_shader_context *ctx)
{
unsigned tcs_out_vertices =
ctx->shader->selector ?
ctx->shader->selector->info.properties[TGSI_PROPERTY_TCS_VERTICES_OUT] : 0;
@@ -440,54 +436,50 @@ static LLVMValueRef get_tcs_in_vertex_dw_stride(struct si_shader_context *ctx)
default:
assert(0);
return NULL;
}
}
static LLVMValueRef get_instance_index_for_fetch(
struct si_shader_context *ctx,
unsigned param_start_instance, LLVMValueRef divisor)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
-
LLVMValueRef result = ctx->abi.instance_id;
/* The division must be done before START_INSTANCE is added. */
if (divisor != ctx->i32_1)
- result = LLVMBuildUDiv(gallivm->builder, result, divisor, "");
+ result = LLVMBuildUDiv(ctx->ac.builder, result, divisor, "");
- return LLVMBuildAdd(gallivm->builder, result,
+ return LLVMBuildAdd(ctx->ac.builder, result,
LLVMGetParam(ctx->main_fn, param_start_instance), "");
}
/* Bitcast <4 x float> to <2 x double>, extract the component, and convert
* to float. */
static LLVMValueRef extract_double_to_float(struct si_shader_context *ctx,
LLVMValueRef vec4,
unsigned double_index)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMTypeRef f64 = LLVMDoubleTypeInContext(ctx->gallivm.context);
LLVMValueRef dvec2 = LLVMBuildBitCast(builder, vec4,
LLVMVectorType(f64, 2), "");
LLVMValueRef index = LLVMConstInt(ctx->i32, double_index, 0);
LLVMValueRef value = LLVMBuildExtractElement(builder, dvec2, index, "");
return LLVMBuildFPTrunc(builder, value, ctx->f32, "");
}
void si_llvm_load_input_vs(
struct si_shader_context *ctx,
unsigned input_index,
LLVMValueRef out[4])
{
- struct gallivm_state *gallivm = &ctx->gallivm;
-
unsigned chan;
unsigned fix_fetch;
unsigned num_fetches;
unsigned fetch_stride;
LLVMValueRef t_list_ptr;
LLVMValueRef t_offset;
LLVMValueRef t_list;
LLVMValueRef vertex_index;
LLVMValueRef input[3];
@@ -534,112 +526,112 @@ void si_llvm_load_input_vs(
LLVMValueRef voffset = LLVMConstInt(ctx->i32, fetch_stride * i, 0);
input[i] = ac_build_buffer_load_format(&ctx->ac, t_list,
vertex_index, voffset,
true);
}
/* Break up the vec4 into individual components */
for (chan = 0; chan < 4; chan++) {
LLVMValueRef llvm_chan = LLVMConstInt(ctx->i32, chan, 0);
- out[chan] = LLVMBuildExtractElement(gallivm->builder,
+ out[chan] = LLVMBuildExtractElement(ctx->ac.builder,
input[0], llvm_chan, "");
}
switch (fix_fetch) {
case SI_FIX_FETCH_A2_SNORM:
case SI_FIX_FETCH_A2_SSCALED:
case SI_FIX_FETCH_A2_SINT: {
/* The hardware returns an unsigned value; convert it to a
* signed one.
*/
LLVMValueRef tmp = out[3];
LLVMValueRef c30 = LLVMConstInt(ctx->i32, 30, 0);
/* First, recover the sign-extended signed integer value. */
if (fix_fetch == SI_FIX_FETCH_A2_SSCALED)
- tmp = LLVMBuildFPToUI(gallivm->builder, tmp, ctx->i32, "");
+ tmp = LLVMBuildFPToUI(ctx->ac.builder, tmp, ctx->i32, "");
else
tmp = ac_to_integer(&ctx->ac, tmp);
/* For the integer-like cases, do a natural sign extension.
*
* For the SNORM case, the values are 0.0, 0.333, 0.666, 1.0
* and happen to contain 0, 1, 2, 3 as the two LSBs of the
* exponent.
*/
- tmp = LLVMBuildShl(gallivm->builder, tmp,
+ tmp = LLVMBuildShl(ctx->ac.builder, tmp,
fix_fetch == SI_FIX_FETCH_A2_SNORM ?
LLVMConstInt(ctx->i32, 7, 0) : c30, "");
- tmp = LLVMBuildAShr(gallivm->builder, tmp, c30, "");
+ tmp = LLVMBuildAShr(ctx->ac.builder, tmp, c30, "");
/* Convert back to the right type. */
if (fix_fetch == SI_FIX_FETCH_A2_SNORM) {
LLVMValueRef clamp;
LLVMValueRef neg_one = LLVMConstReal(ctx->f32, -1.0);
- tmp = LLVMBuildSIToFP(gallivm->builder, tmp, ctx->f32, "");
- clamp = LLVMBuildFCmp(gallivm->builder, LLVMRealULT, tmp, neg_one, "");
- tmp = LLVMBuildSelect(gallivm->builder, clamp, neg_one, tmp, "");
+ tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, "");
+ clamp = LLVMBuildFCmp(ctx->ac.builder, LLVMRealULT, tmp, neg_one, "");
+ tmp = LLVMBuildSelect(ctx->ac.builder, clamp, neg_one, tmp, "");
} else if (fix_fetch == SI_FIX_FETCH_A2_SSCALED) {
- tmp = LLVMBuildSIToFP(gallivm->builder, tmp, ctx->f32, "");
+ tmp = LLVMBuildSIToFP(ctx->ac.builder, tmp, ctx->f32, "");
}
out[3] = tmp;
break;
}
case SI_FIX_FETCH_RGBA_32_UNORM:
case SI_FIX_FETCH_RGBX_32_UNORM:
for (chan = 0; chan < 4; chan++) {
out[chan] = ac_to_integer(&ctx->ac, out[chan]);
- out[chan] = LLVMBuildUIToFP(gallivm->builder,
+ out[chan] = LLVMBuildUIToFP(ctx->ac.builder,
out[chan], ctx->f32, "");
- out[chan] = LLVMBuildFMul(gallivm->builder, out[chan],
+ out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan],
LLVMConstReal(ctx->f32, 1.0 / UINT_MAX), "");
}
/* RGBX UINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
if (fix_fetch == SI_FIX_FETCH_RGBX_32_UNORM)
out[3] = LLVMConstReal(ctx->f32, 1);
break;
case SI_FIX_FETCH_RGBA_32_SNORM:
case SI_FIX_FETCH_RGBX_32_SNORM:
case SI_FIX_FETCH_RGBA_32_FIXED:
case SI_FIX_FETCH_RGBX_32_FIXED: {
double scale;
if (fix_fetch >= SI_FIX_FETCH_RGBA_32_FIXED)
scale = 1.0 / 0x10000;
else
scale = 1.0 / INT_MAX;
for (chan = 0; chan < 4; chan++) {
out[chan] = ac_to_integer(&ctx->ac, out[chan]);
- out[chan] = LLVMBuildSIToFP(gallivm->builder,
+ out[chan] = LLVMBuildSIToFP(ctx->ac.builder,
out[chan], ctx->f32, "");
- out[chan] = LLVMBuildFMul(gallivm->builder, out[chan],
+ out[chan] = LLVMBuildFMul(ctx->ac.builder, out[chan],
LLVMConstReal(ctx->f32, scale), "");
}
/* RGBX SINT returns 1 in alpha, which would be rounded to 0 by normalizing. */
if (fix_fetch == SI_FIX_FETCH_RGBX_32_SNORM ||
fix_fetch == SI_FIX_FETCH_RGBX_32_FIXED)
out[3] = LLVMConstReal(ctx->f32, 1);
break;
}
case SI_FIX_FETCH_RGBA_32_USCALED:
for (chan = 0; chan < 4; chan++) {
out[chan] = ac_to_integer(&ctx->ac, out[chan]);
- out[chan] = LLVMBuildUIToFP(gallivm->builder,
+ out[chan] = LLVMBuildUIToFP(ctx->ac.builder,
out[chan], ctx->f32, "");
}
break;
case SI_FIX_FETCH_RGBA_32_SSCALED:
for (chan = 0; chan < 4; chan++) {
out[chan] = ac_to_integer(&ctx->ac, out[chan]);
- out[chan] = LLVMBuildSIToFP(gallivm->builder,
+ out[chan] = LLVMBuildSIToFP(ctx->ac.builder,
out[chan], ctx->f32, "");
}
break;
case SI_FIX_FETCH_RG_64_FLOAT:
for (chan = 0; chan < 2; chan++)
out[chan] = extract_double_to_float(ctx, input[0], chan);
out[2] = LLVMConstReal(ctx->f32, 0);
out[3] = LLVMConstReal(ctx->f32, 1);
break;
@@ -653,21 +645,21 @@ void si_llvm_load_input_vs(
for (chan = 0; chan < 4; chan++) {
out[chan] = extract_double_to_float(ctx, input[chan / 2],
chan % 2);
}
break;
case SI_FIX_FETCH_RGB_8:
case SI_FIX_FETCH_RGB_8_INT:
case SI_FIX_FETCH_RGB_16:
case SI_FIX_FETCH_RGB_16_INT:
for (chan = 0; chan < 3; chan++) {
- out[chan] = LLVMBuildExtractElement(gallivm->builder,
+ out[chan] = LLVMBuildExtractElement(ctx->ac.builder,
input[chan],
ctx->i32_0, "");
}
if (fix_fetch == SI_FIX_FETCH_RGB_8 ||
fix_fetch == SI_FIX_FETCH_RGB_16) {
out[3] = LLVMConstReal(ctx->f32, 1);
} else {
out[3] = ac_to_float(&ctx->ac, ctx->i32_1);
}
break;
@@ -710,47 +702,45 @@ static LLVMValueRef get_primitive_id(struct si_shader_context *ctx,
/**
* Return the value of tgsi_ind_register for indexing.
* This is the indirect index with the constant offset added to it.
*/
LLVMValueRef si_get_indirect_index(struct si_shader_context *ctx,
const struct tgsi_ind_register *ind,
unsigned addr_mul,
int rel_index)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result;
-
if (ind->File == TGSI_FILE_ADDRESS) {
result = ctx->addrs[ind->Index][ind->Swizzle];
- result = LLVMBuildLoad(gallivm->builder, result, "");
+ result = LLVMBuildLoad(ctx->ac.builder, result, "");
} else {
struct tgsi_full_src_register src = {};
src.Register.File = ind->File;
src.Register.Index = ind->Index;
/* Set the second index to 0 for constants. */
if (ind->File == TGSI_FILE_CONSTANT)
src.Register.Dimension = 1;
result = ctx->bld_base.emit_fetch_funcs[ind->File](&ctx->bld_base, &src,
TGSI_TYPE_SIGNED,
ind->Swizzle);
result = ac_to_integer(&ctx->ac, result);
}
if (addr_mul != 1)
- result = LLVMBuildMul(gallivm->builder, result,
+ result = LLVMBuildMul(ctx->ac.builder, result,
LLVMConstInt(ctx->i32, addr_mul, 0), "");
- result = LLVMBuildAdd(gallivm->builder, result,
+ result = LLVMBuildAdd(ctx->ac.builder, result,
LLVMConstInt(ctx->i32, rel_index, 0), "");
return result;
}
/**
* Like si_get_indirect_index, but restricts the return value to a (possibly
* undefined) value inside [0..num).
*/
LLVMValueRef si_get_bounded_indirect_index(struct si_shader_context *ctx,
const struct tgsi_ind_register *ind,
@@ -764,21 +754,20 @@ LLVMValueRef si_get_bounded_indirect_index(struct si_shader_context *ctx,
/**
* Calculate a dword address given an input or output register and a stride.
*/
static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
const struct tgsi_full_dst_register *dst,
const struct tgsi_full_src_register *src,
LLVMValueRef vertex_dw_stride,
LLVMValueRef base_addr)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
ubyte *name, *index, *array_first;
int first, param;
struct tgsi_full_dst_register reg;
/* Set the register description. The address computation is the same
* for sources and destinations. */
if (src) {
reg.Register.File = src->Register.File;
reg.Register.Index = src->Register.Index;
@@ -794,22 +783,22 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
* in a primitive), calculate the base address of the vertex. */
if (reg.Register.Dimension) {
LLVMValueRef index;
if (reg.Dimension.Indirect)
index = si_get_indirect_index(ctx, ®.DimIndirect,
1, reg.Dimension.Index);
else
index = LLVMConstInt(ctx->i32, reg.Dimension.Index, 0);
- base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
- LLVMBuildMul(gallivm->builder, index,
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMBuildMul(ctx->ac.builder, index,
vertex_dw_stride, ""), "");
}
/* Get information about the register. */
if (reg.Register.File == TGSI_FILE_INPUT) {
name = info->input_semantic_name;
index = info->input_semantic_index;
array_first = info->input_array_first;
} else if (reg.Register.File == TGSI_FILE_OUTPUT) {
name = info->output_semantic_name;
@@ -825,37 +814,37 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
LLVMValueRef ind_index;
if (reg.Indirect.ArrayID)
first = array_first[reg.Indirect.ArrayID];
else
first = reg.Register.Index;
ind_index = si_get_indirect_index(ctx, ®.Indirect,
1, reg.Register.Index - first);
- base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
- LLVMBuildMul(gallivm->builder, ind_index,
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMBuildMul(ctx->ac.builder, ind_index,
LLVMConstInt(ctx->i32, 4, 0), ""), "");
param = reg.Register.Dimension ?
si_shader_io_get_unique_index(name[first], index[first]) :
si_shader_io_get_unique_index_patch(name[first], index[first]);
} else {
param = reg.Register.Dimension ?
si_shader_io_get_unique_index(name[reg.Register.Index],
index[reg.Register.Index]) :
si_shader_io_get_unique_index_patch(name[reg.Register.Index],
index[reg.Register.Index]);
}
/* Add the base address of the element. */
- return LLVMBuildAdd(gallivm->builder, base_addr,
+ return LLVMBuildAdd(ctx->ac.builder, base_addr,
LLVMConstInt(ctx->i32, param * 4, 0), "");
}
/* The offchip buffer layout for TCS->TES is
*
* - attribute 0 of patch 0 vertex 0
* - attribute 0 of patch 0 vertex 1
* - attribute 0 of patch 0 vertex 2
* ...
* - attribute 0 of patch 1 vertex 0
@@ -868,65 +857,63 @@ static LLVMValueRef get_dw_address(struct si_shader_context *ctx,
* - per patch attribute 0 of patch 1
* ...
*
* Note that every attribute has 4 components.
*/
static LLVMValueRef get_tcs_tes_buffer_address(struct si_shader_context *ctx,
LLVMValueRef rel_patch_id,
LLVMValueRef vertex_index,
LLVMValueRef param_index)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef base_addr, vertices_per_patch, num_patches, total_vertices;
LLVMValueRef param_stride, constant16;
vertices_per_patch = get_num_tcs_out_vertices(ctx);
num_patches = unpack_param(ctx, ctx->param_tcs_offchip_layout, 0, 6);
- total_vertices = LLVMBuildMul(gallivm->builder, vertices_per_patch,
+ total_vertices = LLVMBuildMul(ctx->ac.builder, vertices_per_patch,
num_patches, "");
constant16 = LLVMConstInt(ctx->i32, 16, 0);
if (vertex_index) {
- base_addr = LLVMBuildMul(gallivm->builder, rel_patch_id,
+ base_addr = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
vertices_per_patch, "");
- base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
vertex_index, "");
param_stride = total_vertices;
} else {
base_addr = rel_patch_id;
param_stride = num_patches;
}
- base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
- LLVMBuildMul(gallivm->builder, param_index,
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
+ LLVMBuildMul(ctx->ac.builder, param_index,
param_stride, ""), "");
- base_addr = LLVMBuildMul(gallivm->builder, base_addr, constant16, "");
+ base_addr = LLVMBuildMul(ctx->ac.builder, base_addr, constant16, "");
if (!vertex_index) {
LLVMValueRef patch_data_offset =
unpack_param(ctx, ctx->param_tcs_offchip_layout, 12, 20);
- base_addr = LLVMBuildAdd(gallivm->builder, base_addr,
+ base_addr = LLVMBuildAdd(ctx->ac.builder, base_addr,
patch_data_offset, "");
}
return base_addr;
}
static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
struct si_shader_context *ctx,
const struct tgsi_full_dst_register *dst,
const struct tgsi_full_src_register *src)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct tgsi_shader_info *info = &ctx->shader->selector->info;
ubyte *name, *index, *array_first;
struct tgsi_full_src_register reg;
LLVMValueRef vertex_index = NULL;
LLVMValueRef param_index = NULL;
unsigned param_index_base, param_base;
reg = src ? *src : tgsi_full_src_register_from_dst(dst);
if (reg.Register.Dimension) {
@@ -963,52 +950,51 @@ static LLVMValueRef get_tcs_tes_buffer_address_from_reg(
} else {
param_base = reg.Register.Index;
param_index = ctx->i32_0;
}
param_index_base = reg.Register.Dimension ?
si_shader_io_get_unique_index(name[param_base], index[param_base]) :
si_shader_io_get_unique_index_patch(name[param_base], index[param_base]);
- param_index = LLVMBuildAdd(gallivm->builder, param_index,
+ param_index = LLVMBuildAdd(ctx->ac.builder, param_index,
LLVMConstInt(ctx->i32, param_index_base, 0),
"");
return get_tcs_tes_buffer_address(ctx, get_rel_patch_id(ctx),
vertex_index, param_index);
}
static LLVMValueRef buffer_load(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type, unsigned swizzle,
LLVMValueRef buffer, LLVMValueRef offset,
LLVMValueRef base, bool can_speculate)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef value, value2;
LLVMTypeRef llvm_type = tgsi2llvmtype(bld_base, type);
LLVMTypeRef vec_type = LLVMVectorType(llvm_type, 4);
if (swizzle == ~0) {
value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
0, 1, 0, can_speculate, false);
- return LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
+ return LLVMBuildBitCast(ctx->ac.builder, value, vec_type, "");
}
if (!tgsi_type_is_64bit(type)) {
value = ac_build_buffer_load(&ctx->ac, buffer, 4, NULL, base, offset,
0, 1, 0, can_speculate, false);
- value = LLVMBuildBitCast(gallivm->builder, value, vec_type, "");
- return LLVMBuildExtractElement(gallivm->builder, value,
+ value = LLVMBuildBitCast(ctx->ac.builder, value, vec_type, "");
+ return LLVMBuildExtractElement(ctx->ac.builder, value,
LLVMConstInt(ctx->i32, swizzle, 0), "");
}
value = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
swizzle * 4, 1, 0, can_speculate, false);
value2 = ac_build_buffer_load(&ctx->ac, buffer, 1, NULL, base, offset,
swizzle * 4 + 4, 1, 0, can_speculate, false);
return si_llvm_emit_fetch_64bit(bld_base, type, value, value2);
@@ -1071,21 +1057,21 @@ static void lds_store(struct lp_build_tgsi_context *bld_base,
LLVMConstInt(ctx->i32, dw_offset_imm, 0));
value = ac_to_integer(&ctx->ac, value);
ac_build_indexed_store(&ctx->ac, ctx->lds,
dw_addr, value);
}
static LLVMValueRef desc_from_addr_base64k(struct si_shader_context *ctx,
unsigned param)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef addr = LLVMGetParam(ctx->main_fn, param);
addr = LLVMBuildZExt(builder, addr, ctx->i64, "");
addr = LLVMBuildShl(builder, addr, LLVMConstInt(ctx->i64, 16, 0), "");
uint64_t desc2 = 0xffffffff;
uint64_t desc3 = S_008F0C_DST_SEL_X(V_008F0C_SQ_SEL_X) |
S_008F0C_DST_SEL_Y(V_008F0C_SQ_SEL_Y) |
S_008F0C_DST_SEL_Z(V_008F0C_SQ_SEL_Z) |
S_008F0C_DST_SEL_W(V_008F0C_SQ_SEL_W) |
@@ -1224,24 +1210,24 @@ static void store_output_tcs(struct lp_build_tgsi_context *bld_base,
if (reg->Register.WriteMask != 0xF && !is_tess_factor) {
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 1,
buf_addr, base,
4 * chan_index, 1, 0, true, false);
}
/* Write tess factors into VGPRs for the epilog. */
if (is_tess_factor &&
ctx->shader->selector->tcs_info.tessfactors_are_def_in_all_invocs) {
if (!is_tess_inner) {
- LLVMBuildStore(gallivm->builder, value, /* outer */
+ LLVMBuildStore(ctx->ac.builder, value, /* outer */
ctx->invoc0_tess_factors[chan_index]);
} else if (chan_index < 2) {
- LLVMBuildStore(gallivm->builder, value, /* inner */
+ LLVMBuildStore(ctx->ac.builder, value, /* inner */
ctx->invoc0_tess_factors[4 + chan_index]);
}
}
}
if (reg->Register.WriteMask == 0xF && !is_tess_factor) {
LLVMValueRef value = lp_build_gather_values(gallivm,
values, 4);
ac_build_buffer_store_dword(&ctx->ac, buffer, value, 4, buf_addr,
base, 0, 1, 0, true, false);
@@ -1288,21 +1274,21 @@ static LLVMValueRef fetch_input_gs(
break;
case 2:
vtx_offset = unpack_param(ctx, ctx->param_gs_vtx45_offset,
index % 2 ? 16 : 0, 16);
break;
default:
assert(0);
return NULL;
}
- vtx_offset = LLVMBuildAdd(gallivm->builder, vtx_offset,
+ vtx_offset = LLVMBuildAdd(ctx->ac.builder, vtx_offset,
LLVMConstInt(ctx->i32, param * 4, 0), "");
return lds_load(bld_base, type, swizzle, vtx_offset);
}
/* GFX6: input load from the ESGS ring in memory. */
if (swizzle == ~0) {
LLVMValueRef values[TGSI_NUM_CHANNELS];
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
values[chan] = fetch_input_gs(bld_base, reg, type, chan);
@@ -1406,74 +1392,73 @@ static void interp_fs_input(struct si_shader_context *ctx,
unsigned input_index,
unsigned semantic_name,
unsigned semantic_index,
unsigned num_interp_inputs,
unsigned colors_read_mask,
LLVMValueRef interp_param,
LLVMValueRef prim_mask,
LLVMValueRef face,
LLVMValueRef result[4])
{
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef i = NULL, j = NULL;
unsigned chan;
/* fs.constant returns the param from the middle vertex, so it's not
* really useful for flat shading. It's meant to be used for custom
* interpolation (but the intrinsic can't fetch from the other two
* vertices).
*
* Luckily, it doesn't matter, because we rely on the FLAT_SHADE state
* to do the right thing. The only reason we use fs.constant is that
* fs.interp cannot be used on integers, because they can be equal
* to NaN.
*
* When interp is false we will use fs.constant or for newer llvm,
* amdgcn.interp.mov.
*/
bool interp = interp_param != NULL;
if (interp) {
- interp_param = LLVMBuildBitCast(gallivm->builder, interp_param,
+ interp_param = LLVMBuildBitCast(ctx->ac.builder, interp_param,
LLVMVectorType(ctx->f32, 2), "");
- i = LLVMBuildExtractElement(gallivm->builder, interp_param,
+ i = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
ctx->i32_0, "");
- j = LLVMBuildExtractElement(gallivm->builder, interp_param,
+ j = LLVMBuildExtractElement(ctx->ac.builder, interp_param,
ctx->i32_1, "");
}
if (semantic_name == TGSI_SEMANTIC_COLOR &&
ctx->shader->key.part.ps.prolog.color_two_side) {
LLVMValueRef is_face_positive;
/* If BCOLOR0 is used, BCOLOR1 is at offset "num_inputs + 1",
* otherwise it's at offset "num_inputs".
*/
unsigned back_attr_offset = num_interp_inputs;
if (semantic_index == 1 && colors_read_mask & 0xf)
back_attr_offset += 1;
- is_face_positive = LLVMBuildICmp(gallivm->builder, LLVMIntNE,
+ is_face_positive = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
face, ctx->i32_0, "");
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
LLVMValueRef front, back;
front = si_build_fs_interp(ctx,
input_index, chan,
prim_mask, i, j);
back = si_build_fs_interp(ctx,
back_attr_offset, chan,
prim_mask, i, j);
- result[chan] = LLVMBuildSelect(gallivm->builder,
+ result[chan] = LLVMBuildSelect(ctx->ac.builder,
is_face_positive,
front,
back,
"");
}
} else if (semantic_name == TGSI_SEMANTIC_FOG) {
result[0] = si_build_fs_interp(ctx, input_index,
0, prim_mask, i, j);
result[1] =
result[2] = LLVMConstReal(ctx->f32, 0.0f);
@@ -1555,28 +1540,27 @@ static LLVMValueRef buffer_load_const(struct si_shader_context *ctx,
LLVMValueRef offset)
{
return ac_build_buffer_load(&ctx->ac, resource, 1, NULL, offset, NULL,
0, 0, 0, true, true);
}
static LLVMValueRef load_sample_position(struct si_shader_context *ctx, LLVMValueRef sample_id)
{
struct lp_build_context *uint_bld = &ctx->bld_base.uint_bld;
struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef desc = LLVMGetParam(ctx->main_fn, ctx->param_rw_buffers);
LLVMValueRef buf_index = LLVMConstInt(ctx->i32, SI_PS_CONST_SAMPLE_POSITIONS, 0);
LLVMValueRef resource = ac_build_indexed_load_const(&ctx->ac, desc, buf_index);
/* offset = sample_id * 8 (8 = 2 floats containing samplepos.xy) */
LLVMValueRef offset0 = lp_build_mul_imm(uint_bld, sample_id, 8);
- LLVMValueRef offset1 = LLVMBuildAdd(builder, offset0, LLVMConstInt(ctx->i32, 4, 0), "");
+ LLVMValueRef offset1 = LLVMBuildAdd(ctx->ac.builder, offset0, LLVMConstInt(ctx->i32, 4, 0), "");
LLVMValueRef pos[4] = {
buffer_load_const(ctx, resource, offset0),
buffer_load_const(ctx, resource, offset1),
LLVMConstReal(ctx->f32, 0),
LLVMConstReal(ctx->f32, 0)
};
return lp_build_gather_values(gallivm, pos, 4);
}
@@ -1590,44 +1574,44 @@ void si_load_system_value(struct si_shader_context *ctx,
LLVMValueRef value = 0;
assert(index < RADEON_LLVM_MAX_SYSTEM_VALUES);
switch (decl->Semantic.Name) {
case TGSI_SEMANTIC_INSTANCEID:
value = ctx->abi.instance_id;
break;
case TGSI_SEMANTIC_VERTEXID:
- value = LLVMBuildAdd(gallivm->builder,
+ value = LLVMBuildAdd(ctx->ac.builder,
ctx->abi.vertex_id,
ctx->abi.base_vertex, "");
break;
case TGSI_SEMANTIC_VERTEXID_NOBASE:
/* Unused. Clarify the meaning in indexed vs. non-indexed
* draws if this is ever used again. */
assert(false);
break;
case TGSI_SEMANTIC_BASEVERTEX:
{
/* For non-indexed draws, the base vertex set by the driver
* (for direct draws) or the CP (for indirect draws) is the
* first vertex ID, but GLSL expects 0 to be returned.
*/
LLVMValueRef vs_state = LLVMGetParam(ctx->main_fn, ctx->param_vs_state_bits);
LLVMValueRef indexed;
- indexed = LLVMBuildLShr(gallivm->builder, vs_state, ctx->i32_1, "");
- indexed = LLVMBuildTrunc(gallivm->builder, indexed, ctx->i1, "");
+ indexed = LLVMBuildLShr(ctx->ac.builder, vs_state, ctx->i32_1, "");
+ indexed = LLVMBuildTrunc(ctx->ac.builder, indexed, ctx->i1, "");
- value = LLVMBuildSelect(gallivm->builder, indexed,
+ value = LLVMBuildSelect(ctx->ac.builder, indexed,
ctx->abi.base_vertex, ctx->i32_0, "");
break;
}
case TGSI_SEMANTIC_BASEINSTANCE:
value = ctx->abi.start_instance;
break;
case TGSI_SEMANTIC_DRAWID:
value = ctx->abi.draw_id;
@@ -1795,65 +1779,65 @@ void si_load_system_value(struct si_shader_context *ctx,
}
value = lp_build_gather_values(gallivm, values, 3);
break;
}
case TGSI_SEMANTIC_THREAD_ID:
value = LLVMGetParam(ctx->main_fn, ctx->param_thread_id);
break;
case TGSI_SEMANTIC_HELPER_INVOCATION:
- value = lp_build_intrinsic(gallivm->builder,
+ value = lp_build_intrinsic(ctx->ac.builder,
"llvm.amdgcn.ps.live",
ctx->i1, NULL, 0,
LP_FUNC_ATTR_READNONE);
- value = LLVMBuildNot(gallivm->builder, value, "");
- value = LLVMBuildSExt(gallivm->builder, value, ctx->i32, "");
+ value = LLVMBuildNot(ctx->ac.builder, value, "");
+ value = LLVMBuildSExt(ctx->ac.builder, value, ctx->i32, "");
break;
case TGSI_SEMANTIC_SUBGROUP_SIZE:
value = LLVMConstInt(ctx->i32, 64, 0);
break;
case TGSI_SEMANTIC_SUBGROUP_INVOCATION:
value = ac_get_thread_id(&ctx->ac);
break;
case TGSI_SEMANTIC_SUBGROUP_EQ_MASK:
{
LLVMValueRef id = ac_get_thread_id(&ctx->ac);
- id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, "");
- value = LLVMBuildShl(gallivm->builder, LLVMConstInt(ctx->i64, 1, 0), id, "");
- value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, "");
+ id = LLVMBuildZExt(ctx->ac.builder, id, ctx->i64, "");
+ value = LLVMBuildShl(ctx->ac.builder, LLVMConstInt(ctx->i64, 1, 0), id, "");
+ value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->v2i32, "");
break;
}
case TGSI_SEMANTIC_SUBGROUP_GE_MASK:
case TGSI_SEMANTIC_SUBGROUP_GT_MASK:
case TGSI_SEMANTIC_SUBGROUP_LE_MASK:
case TGSI_SEMANTIC_SUBGROUP_LT_MASK:
{
LLVMValueRef id = ac_get_thread_id(&ctx->ac);
if (decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_GT_MASK ||
decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LE_MASK) {
/* All bits set except LSB */
value = LLVMConstInt(ctx->i64, -2, 0);
} else {
/* All bits set */
value = LLVMConstInt(ctx->i64, -1, 0);
}
- id = LLVMBuildZExt(gallivm->builder, id, ctx->i64, "");
- value = LLVMBuildShl(gallivm->builder, value, id, "");
+ id = LLVMBuildZExt(ctx->ac.builder, id, ctx->i64, "");
+ value = LLVMBuildShl(ctx->ac.builder, value, id, "");
if (decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LE_MASK ||
decl->Semantic.Name == TGSI_SEMANTIC_SUBGROUP_LT_MASK)
- value = LLVMBuildNot(gallivm->builder, value, "");
- value = LLVMBuildBitCast(gallivm->builder, value, ctx->v2i32, "");
+ value = LLVMBuildNot(ctx->ac.builder, value, "");
+ value = LLVMBuildBitCast(ctx->ac.builder, value, ctx->v2i32, "");
break;
}
default:
assert(!"unknown system value");
return;
}
ctx->system_values[index] = value;
}
@@ -1870,53 +1854,53 @@ void si_declare_compute_memory(struct si_shader_context *ctx,
assert(decl->Declaration.MemType == TGSI_MEMORY_TYPE_SHARED);
assert(decl->Range.First == decl->Range.Last);
assert(!ctx->shared_memory);
var = LLVMAddGlobalInAddressSpace(gallivm->module,
LLVMArrayType(ctx->i8, sel->local_size),
"compute_lds",
LOCAL_ADDR_SPACE);
LLVMSetAlignment(var, 4);
- ctx->shared_memory = LLVMBuildBitCast(gallivm->builder, var, i8p, "");
+ ctx->shared_memory = LLVMBuildBitCast(ctx->ac.builder, var, i8p, "");
}
static LLVMValueRef load_const_buffer_desc(struct si_shader_context *ctx, int i)
{
LLVMValueRef list_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_const_and_shader_buffers);
return ac_build_indexed_load_const(&ctx->ac, list_ptr,
LLVMConstInt(ctx->i32, si_get_constbuf_slot(i), 0));
}
static LLVMValueRef load_ubo(struct ac_shader_abi *abi, LLVMValueRef index)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers);
index = si_llvm_bound_index(ctx, index, ctx->num_const_buffers);
- index = LLVMBuildAdd(ctx->gallivm.builder, index,
+ index = LLVMBuildAdd(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), "");
return ac_build_indexed_load_const(&ctx->ac, ptr, index);
}
static LLVMValueRef
load_ssbo(struct ac_shader_abi *abi, LLVMValueRef index, bool write)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
LLVMValueRef rsrc_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_const_and_shader_buffers);
index = si_llvm_bound_index(ctx, index, ctx->num_shader_buffers);
- index = LLVMBuildSub(ctx->gallivm.builder,
+ index = LLVMBuildSub(ctx->ac.builder,
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS - 1, 0),
index, "");
return ac_build_indexed_load_const(&ctx->ac, rsrc_ptr, index);
}
static LLVMValueRef fetch_constant(
struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
@@ -1941,21 +1925,21 @@ static LLVMValueRef fetch_constant(
assert(reg->Register.Dimension);
buf = reg->Dimension.Index;
idx = reg->Register.Index * 4 + swizzle;
if (reg->Dimension.Indirect) {
LLVMValueRef ptr = LLVMGetParam(ctx->main_fn, ctx->param_const_and_shader_buffers);
LLVMValueRef index;
index = si_get_bounded_indirect_index(ctx, ®->DimIndirect,
reg->Dimension.Index,
ctx->num_const_buffers);
- index = LLVMBuildAdd(ctx->gallivm.builder, index,
+ index = LLVMBuildAdd(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, SI_NUM_SHADER_BUFFERS, 0), "");
bufp = ac_build_indexed_load_const(&ctx->ac, ptr, index);
} else
bufp = load_const_buffer_desc(ctx, buf);
if (reg->Register.Indirect) {
addr = si_get_indirect_index(ctx, ireg, 16, idx * 4);
} else {
addr = LLVMConstInt(ctx->i32, idx * 4, 0);
}
@@ -1974,47 +1958,47 @@ static LLVMValueRef fetch_constant(
result = si_llvm_emit_fetch_64bit(bld_base, type,
result, result2);
}
return result;
}
/* Upper 16 bits must be zero. */
static LLVMValueRef si_llvm_pack_two_int16(struct si_shader_context *ctx,
LLVMValueRef val[2])
{
- return LLVMBuildOr(ctx->gallivm.builder, val[0],
- LLVMBuildShl(ctx->gallivm.builder, val[1],
+ return LLVMBuildOr(ctx->ac.builder, val[0],
+ LLVMBuildShl(ctx->ac.builder, val[1],
LLVMConstInt(ctx->i32, 16, 0),
""), "");
}
/* Upper 16 bits are ignored and will be dropped. */
static LLVMValueRef si_llvm_pack_two_int32_as_int16(struct si_shader_context *ctx,
LLVMValueRef val[2])
{
LLVMValueRef v[2] = {
- LLVMBuildAnd(ctx->gallivm.builder, val[0],
+ LLVMBuildAnd(ctx->ac.builder, val[0],
LLVMConstInt(ctx->i32, 0xffff, 0), ""),
val[1],
};
return si_llvm_pack_two_int16(ctx, v);
}
/* Initialize arguments for the shader export intrinsic */
static void si_llvm_init_export_args(struct lp_build_tgsi_context *bld_base,
LLVMValueRef *values,
unsigned target,
struct ac_export_args *args)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct lp_build_context *base = &bld_base->base;
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef val[4];
unsigned spi_shader_col_format = V_028714_SPI_SHADER_32_ABGR;
unsigned chan;
bool is_int8, is_int10;
/* Default is 0xf. Adjusted below depending on the format. */
args->enabled_channels = 0xf; /* writemask */
/* Specify whether the EXEC mask represents the valid mask */
args->valid_mask = 0;
@@ -2199,40 +2183,39 @@ static void si_alpha_test(struct lp_build_tgsi_context *bld_base,
} else {
ac_build_kill(&ctx->ac, NULL);
}
}
static LLVMValueRef si_scale_alpha_by_sample_mask(struct lp_build_tgsi_context *bld_base,
LLVMValueRef alpha,
unsigned samplemask_param)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef coverage;
/* alpha = alpha * popcount(coverage) / SI_NUM_SMOOTH_AA_SAMPLES */
coverage = LLVMGetParam(ctx->main_fn,
samplemask_param);
coverage = ac_to_integer(&ctx->ac, coverage);
- coverage = lp_build_intrinsic(gallivm->builder, "llvm.ctpop.i32",
+ coverage = lp_build_intrinsic(ctx->ac.builder, "llvm.ctpop.i32",
ctx->i32,
&coverage, 1, LP_FUNC_ATTR_READNONE);
- coverage = LLVMBuildUIToFP(gallivm->builder, coverage,
+ coverage = LLVMBuildUIToFP(ctx->ac.builder, coverage,
ctx->f32, "");
- coverage = LLVMBuildFMul(gallivm->builder, coverage,
+ coverage = LLVMBuildFMul(ctx->ac.builder, coverage,
LLVMConstReal(ctx->f32,
1.0 / SI_NUM_SMOOTH_AA_SAMPLES), "");
- return LLVMBuildFMul(gallivm->builder, alpha, coverage, "");
+ return LLVMBuildFMul(ctx->ac.builder, alpha, coverage, "");
}
static void si_llvm_emit_clipvertex(struct lp_build_tgsi_context *bld_base,
struct ac_export_args *pos, LLVMValueRef *out_elts)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct lp_build_context *base = &bld_base->base;
unsigned reg_index;
unsigned chan;
unsigned const_chan;
@@ -2293,22 +2276,20 @@ static void si_dump_streamout(struct pipe_stream_output_info *so)
mask & 8 ? "w" : "");
}
}
static void emit_streamout_output(struct si_shader_context *ctx,
LLVMValueRef const *so_buffers,
LLVMValueRef const *so_write_offsets,
struct pipe_stream_output *stream_out,
struct si_shader_output_values *shader_out)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
unsigned buf_idx = stream_out->output_buffer;
unsigned start = stream_out->start_component;
unsigned num_comps = stream_out->num_components;
LLVMValueRef out[4];
assert(num_comps && num_comps <= 4);
if (!num_comps || num_comps > 4)
return;
/* Load the output as int. */
@@ -2323,21 +2304,21 @@ static void emit_streamout_output(struct si_shader_context *ctx,
switch (num_comps) {
case 1: /* as i32 */
vdata = out[0];
break;
case 2: /* as v2i32 */
case 3: /* as v4i32 (aligned to 4) */
case 4: /* as v4i32 */
vdata = LLVMGetUndef(LLVMVectorType(ctx->i32, util_next_power_of_two(num_comps)));
for (int j = 0; j < num_comps; j++) {
- vdata = LLVMBuildInsertElement(builder, vdata, out[j],
+ vdata = LLVMBuildInsertElement(ctx->ac.builder, vdata, out[j],
LLVMConstInt(ctx->i32, j, 0), "");
}
break;
}
ac_build_buffer_store_dword(&ctx->ac, so_buffers[buf_idx],
vdata, num_comps,
so_write_offsets[buf_idx],
ctx->i32_0,
stream_out->dst_offset * 4, 1, 1, true, false);
@@ -2347,21 +2328,21 @@ static void emit_streamout_output(struct si_shader_context *ctx,
* Write streamout data to buffers for vertex stream @p stream (different
* vertex streams can occur for GS copy shaders).
*/
static void si_llvm_emit_streamout(struct si_shader_context *ctx,
struct si_shader_output_values *outputs,
unsigned noutput, unsigned stream)
{
struct si_shader_selector *sel = ctx->shader->selector;
struct pipe_stream_output_info *so = &sel->so;
struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
int i;
struct lp_build_if_state if_ctx;
/* Get bits [22:16], i.e. (so_param >> 16) & 127; */
LLVMValueRef so_vtx_count =
unpack_param(ctx, ctx->param_streamout_config, 16, 7);
LLVMValueRef tid = ac_get_thread_id(&ctx->ac);
/* can_emit = tid < so_vtx_count; */
@@ -2564,21 +2545,21 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
pos_args[1].out[1] = base->zero; /* Y */
pos_args[1].out[2] = base->zero; /* Z */
pos_args[1].out[3] = base->zero; /* W */
if (shader->selector->info.writes_psize)
pos_args[1].out[0] = psize_value;
if (shader->selector->info.writes_edgeflag) {
/* The output is a float, but the hw expects an integer
* with the first bit containing the edge flag. */
- edgeflag_value = LLVMBuildFPToUI(ctx->gallivm.builder,
+ edgeflag_value = LLVMBuildFPToUI(ctx->ac.builder,
edgeflag_value,
ctx->i32, "");
edgeflag_value = ac_build_umin(&ctx->ac,
edgeflag_value,
ctx->i32_1);
/* The LLVM intrinsic expects a float. */
pos_args[1].out[1] = ac_to_float(&ctx->ac, edgeflag_value);
}
@@ -2586,23 +2567,23 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
/* GFX9 has the layer in out.z[10:0] and the viewport
* index in out.z[19:16].
*/
if (shader->selector->info.writes_layer)
pos_args[1].out[2] = layer_value;
if (shader->selector->info.writes_viewport_index) {
LLVMValueRef v = viewport_index_value;
v = ac_to_integer(&ctx->ac, v);
- v = LLVMBuildShl(ctx->gallivm.builder, v,
+ v = LLVMBuildShl(ctx->ac.builder, v,
LLVMConstInt(ctx->i32, 16, 0), "");
- v = LLVMBuildOr(ctx->gallivm.builder, v,
+ v = LLVMBuildOr(ctx->ac.builder, v,
ac_to_integer(&ctx->ac, pos_args[1].out[2]), "");
pos_args[1].out[2] = ac_to_float(&ctx->ac, v);
pos_args[1].enabled_channels |= 1 << 2;
}
} else {
if (shader->selector->info.writes_layer)
pos_args[1].out[2] = layer_value;
if (shader->selector->info.writes_viewport_index) {
pos_args[1].out[3] = viewport_index_value;
@@ -2634,40 +2615,39 @@ static void si_llvm_export_vs(struct lp_build_tgsi_context *bld_base,
si_build_param_exports(ctx, outputs, noutput);
}
/**
* Forward all outputs from the vertex shader to the TES. This is only used
* for the fixed function TCS.
*/
static void si_copy_tcs_inputs(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef invocation_id, buffer, buffer_offset;
LLVMValueRef lds_vertex_stride, lds_vertex_offset, lds_base;
uint64_t inputs;
invocation_id = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_offchip_addr_base64k);
buffer_offset = LLVMGetParam(ctx->main_fn, ctx->param_tcs_offchip_offset);
lds_vertex_stride = get_tcs_in_vertex_dw_stride(ctx);
- lds_vertex_offset = LLVMBuildMul(gallivm->builder, invocation_id,
+ lds_vertex_offset = LLVMBuildMul(ctx->ac.builder, invocation_id,
lds_vertex_stride, "");
lds_base = get_tcs_in_current_patch_offset(ctx);
- lds_base = LLVMBuildAdd(gallivm->builder, lds_base, lds_vertex_offset, "");
+ lds_base = LLVMBuildAdd(ctx->ac.builder, lds_base, lds_vertex_offset, "");
inputs = ctx->shader->key.mono.u.ff_tcs_inputs_to_copy;
while (inputs) {
unsigned i = u_bit_scan64(&inputs);
- LLVMValueRef lds_ptr = LLVMBuildAdd(gallivm->builder, lds_base,
+ LLVMValueRef lds_ptr = LLVMBuildAdd(ctx->ac.builder, lds_base,
LLVMConstInt(ctx->i32, 4 * i, 0),
"");
LLVMValueRef buffer_addr = get_tcs_tes_buffer_address(ctx,
get_rel_patch_id(ctx),
invocation_id,
LLVMConstInt(ctx->i32, i, 0));
LLVMValueRef value = lds_load(bld_base, TGSI_TYPE_SIGNED, ~0,
lds_ptr);
@@ -2697,21 +2677,21 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
if (!shader->key.part.tcs.epilog.invoc0_tess_factors_are_def)
si_llvm_emit_barrier(NULL, bld_base, NULL);
/* Do this only for invocation 0, because the tess levels are per-patch,
* not per-vertex.
*
* This can't jump, because invocation 0 executes this. It should
* at least mask out the loads and stores for other invocations.
*/
lp_build_if(&if_ctx, gallivm,
- LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
+ LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
invocation_id, ctx->i32_0, ""));
/* Determine the layout of one tess factor element in the buffer. */
switch (shader->key.part.tcs.epilog.prim_mode) {
case PIPE_PRIM_LINES:
stride = 2; /* 2 dwords, 1 vec2 store */
outer_comps = 2;
inner_comps = 0;
break;
case PIPE_PRIM_TRIANGLES:
@@ -2741,24 +2721,24 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
for (i = 0; i < inner_comps; i++)
inner[i] = out[outer_comps+i] = invoc0_tf_inner[i];
} else {
/* Load tess_inner and tess_outer from LDS.
* Any invocation can write them, so we can't get them from a temporary.
*/
tess_inner_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSINNER, 0);
tess_outer_index = si_shader_io_get_unique_index_patch(TGSI_SEMANTIC_TESSOUTER, 0);
lds_base = tcs_out_current_patch_data_offset;
- lds_inner = LLVMBuildAdd(gallivm->builder, lds_base,
+ lds_inner = LLVMBuildAdd(ctx->ac.builder, lds_base,
LLVMConstInt(ctx->i32,
tess_inner_index * 4, 0), "");
- lds_outer = LLVMBuildAdd(gallivm->builder, lds_base,
+ lds_outer = LLVMBuildAdd(ctx->ac.builder, lds_base,
LLVMConstInt(ctx->i32,
tess_outer_index * 4, 0), "");
for (i = 0; i < outer_comps; i++) {
outer[i] = out[i] =
lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_outer);
}
for (i = 0; i < inner_comps; i++) {
inner[i] = out[outer_comps+i] =
lds_load(bld_base, TGSI_TYPE_SIGNED, i, lds_inner);
@@ -2780,25 +2760,25 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
if (stride > 4)
vec1 = lp_build_gather_values(gallivm, out+4, stride - 4);
/* Get the buffer. */
buffer = desc_from_addr_base64k(ctx, ctx->param_tcs_factor_addr_base64k);
/* Get the offset. */
tf_base = LLVMGetParam(ctx->main_fn,
ctx->param_tcs_factor_offset);
- byteoffset = LLVMBuildMul(gallivm->builder, rel_patch_id,
+ byteoffset = LLVMBuildMul(ctx->ac.builder, rel_patch_id,
LLVMConstInt(ctx->i32, 4 * stride, 0), "");
lp_build_if(&inner_if_ctx, gallivm,
- LLVMBuildICmp(gallivm->builder, LLVMIntEQ,
+ LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ,
rel_patch_id, ctx->i32_0, ""));
/* Store the dynamic HS control word. */
offset = 0;
if (ctx->screen->b.chip_class <= VI) {
ac_build_buffer_store_dword(&ctx->ac, buffer,
LLVMConstInt(ctx->i32, 0x80000000, 0),
1, ctx->i32_0, tf_base,
offset, 1, 0, true, false);
offset += 4;
@@ -2850,58 +2830,58 @@ static void si_write_tess_factors(struct lp_build_tgsi_context *bld_base,
}
}
lp_build_endif(&if_ctx);
}
static LLVMValueRef
si_insert_input_ret(struct si_shader_context *ctx, LLVMValueRef ret,
unsigned param, unsigned return_index)
{
- return LLVMBuildInsertValue(ctx->gallivm.builder, ret,
+ return LLVMBuildInsertValue(ctx->ac.builder, ret,
LLVMGetParam(ctx->main_fn, param),
return_index, "");
}
static LLVMValueRef
si_insert_input_ret_float(struct si_shader_context *ctx, LLVMValueRef ret,
unsigned param, unsigned return_index)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef p = LLVMGetParam(ctx->main_fn, param);
return LLVMBuildInsertValue(builder, ret,
ac_to_float(&ctx->ac, p),
return_index, "");
}
static LLVMValueRef
si_insert_input_ptr_as_2xi32(struct si_shader_context *ctx, LLVMValueRef ret,
unsigned param, unsigned return_index)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef ptr, lo, hi;
ptr = LLVMGetParam(ctx->main_fn, param);
ptr = LLVMBuildPtrToInt(builder, ptr, ctx->i64, "");
ptr = LLVMBuildBitCast(builder, ptr, ctx->v2i32, "");
lo = LLVMBuildExtractElement(builder, ptr, ctx->i32_0, "");
hi = LLVMBuildExtractElement(builder, ptr, ctx->i32_1, "");
ret = LLVMBuildInsertValue(builder, ret, lo, return_index, "");
return LLVMBuildInsertValue(builder, ret, hi, return_index + 1, "");
}
/* This only writes the tessellation factor levels. */
static void si_llvm_emit_tcs_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef rel_patch_id, invocation_id, tf_lds_offset;
si_copy_tcs_inputs(bld_base);
rel_patch_id = get_rel_patch_id(ctx);
invocation_id = unpack_param(ctx, ctx->param_tcs_rel_ids, 8, 5);
tf_lds_offset = get_tcs_out_current_patch_data_offset(ctx);
if (ctx->screen->b.chip_class >= GFX9) {
LLVMBasicBlockRef blocks[2] = {
@@ -3050,26 +3030,25 @@ static void si_set_es_return_value_for_gs(struct si_shader_context *ctx)
ret = si_insert_input_ret_float(ctx, ret, param, vgpr++);
}
ctx->return_value = ret;
}
static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
- struct gallivm_state *gallivm = &ctx->gallivm;
unsigned i, chan;
LLVMValueRef vertex_id = LLVMGetParam(ctx->main_fn,
ctx->param_rel_auto_id);
LLVMValueRef vertex_dw_stride = get_tcs_in_vertex_dw_stride(ctx);
- LLVMValueRef base_dw_addr = LLVMBuildMul(gallivm->builder, vertex_id,
+ LLVMValueRef base_dw_addr = LLVMBuildMul(ctx->ac.builder, vertex_id,
vertex_dw_stride, "");
/* Write outputs to LDS. The next shader (TCS aka HS) will read
* its inputs from it. */
for (i = 0; i < info->num_outputs; i++) {
LLVMValueRef *out_ptr = ctx->outputs[i];
unsigned name = info->output_semantic_name[i];
unsigned index = info->output_semantic_index[i];
/* The ARB_shader_viewport_layer_array spec contains the
@@ -3085,69 +3064,68 @@ static void si_llvm_emit_ls_epilogue(struct lp_build_tgsi_context *bld_base)
* statically assign to gl_ViewportIndex or gl_Layer, index
* or layer zero is assumed.
*
* So writes to those outputs in VS-as-LS are simply ignored.
*/
if (name == TGSI_SEMANTIC_LAYER ||
name == TGSI_SEMANTIC_VIEWPORT_INDEX)
continue;
int param = si_shader_io_get_unique_index(name, index);
- LLVMValueRef dw_addr = LLVMBuildAdd(gallivm->builder, base_dw_addr,
+ LLVMValueRef dw_addr = LLVMBuildAdd(ctx->ac.builder, base_dw_addr,
LLVMConstInt(ctx->i32, param * 4, 0), "");
for (chan = 0; chan < 4; chan++) {
lds_store(bld_base, chan, dw_addr,
- LLVMBuildLoad(gallivm->builder, out_ptr[chan], ""));
+ LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], ""));
}
}
if (ctx->screen->b.chip_class >= GFX9)
si_set_ls_return_value_for_tcs(ctx);
}
static void si_llvm_emit_es_epilogue(struct lp_build_tgsi_context *bld_base)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_shader *es = ctx->shader;
struct tgsi_shader_info *info = &es->selector->info;
LLVMValueRef soffset = LLVMGetParam(ctx->main_fn,
ctx->param_es2gs_offset);
LLVMValueRef lds_base = NULL;
unsigned chan;
int i;
if (ctx->screen->b.chip_class >= GFX9 && info->num_outputs) {
unsigned itemsize_dw = es->selector->esgs_itemsize / 4;
LLVMValueRef vertex_idx = ac_get_thread_id(&ctx->ac);
LLVMValueRef wave_idx = unpack_param(ctx, ctx->param_merged_wave_info, 24, 4);
- vertex_idx = LLVMBuildOr(gallivm->builder, vertex_idx,
- LLVMBuildMul(gallivm->builder, wave_idx,
+ vertex_idx = LLVMBuildOr(ctx->ac.builder, vertex_idx,
+ LLVMBuildMul(ctx->ac.builder, wave_idx,
LLVMConstInt(ctx->i32, 64, false), ""), "");
- lds_base = LLVMBuildMul(gallivm->builder, vertex_idx,
+ lds_base = LLVMBuildMul(ctx->ac.builder, vertex_idx,
LLVMConstInt(ctx->i32, itemsize_dw, 0), "");
}
for (i = 0; i < info->num_outputs; i++) {
LLVMValueRef *out_ptr = ctx->outputs[i];
int param;
if (info->output_semantic_name[i] == TGSI_SEMANTIC_VIEWPORT_INDEX ||
info->output_semantic_name[i] == TGSI_SEMANTIC_LAYER)
continue;
param = si_shader_io_get_unique_index(info->output_semantic_name[i],
info->output_semantic_index[i]);
for (chan = 0; chan < 4; chan++) {
- LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
+ LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], "");
out_val = ac_to_integer(&ctx->ac, out_val);
/* GFX9 has the ESGS ring in LDS. */
if (ctx->screen->b.chip_class >= GFX9) {
lds_store(bld_base, param * 4 + chan, lds_base, out_val);
continue;
}
ac_build_buffer_store_dword(&ctx->ac,
ctx->esgs_ring,
@@ -3209,44 +3187,44 @@ static void si_llvm_emit_vs_epilogue(struct ac_shader_abi *abi,
for (i = 0; i < info->num_outputs; i++) {
if (info->output_semantic_name[i] != TGSI_SEMANTIC_COLOR &&
info->output_semantic_name[i] != TGSI_SEMANTIC_BCOLOR)
continue;
/* We've found a color. */
if (!cond) {
/* The state is in the first bit of the user SGPR. */
cond = LLVMGetParam(ctx->main_fn,
ctx->param_vs_state_bits);
- cond = LLVMBuildTrunc(gallivm->builder, cond,
+ cond = LLVMBuildTrunc(ctx->ac.builder, cond,
ctx->i1, "");
lp_build_if(&if_ctx, gallivm, cond);
}
for (j = 0; j < 4; j++) {
addr = addrs[4 * i + j];
- val = LLVMBuildLoad(gallivm->builder, addr, "");
+ val = LLVMBuildLoad(ctx->ac.builder, addr, "");
val = ac_build_clamp(&ctx->ac, val);
- LLVMBuildStore(gallivm->builder, val, addr);
+ LLVMBuildStore(ctx->ac.builder, val, addr);
}
}
if (cond)
lp_build_endif(&if_ctx);
}
for (i = 0; i < info->num_outputs; i++) {
outputs[i].semantic_name = info->output_semantic_name[i];
outputs[i].semantic_index = info->output_semantic_index[i];
for (j = 0; j < 4; j++) {
outputs[i].values[j] =
- LLVMBuildLoad(gallivm->builder,
+ LLVMBuildLoad(ctx->ac.builder,
addrs[4 * i + j],
"");
outputs[i].vertex_stream[j] =
(info->output_streams[i] >> (2 * j)) & 3;
}
}
if (ctx->shader->selector->so.num_outputs)
si_llvm_emit_streamout(ctx, outputs, i, 0);
@@ -3325,21 +3303,21 @@ static void si_export_mrt_z(struct lp_build_tgsi_context *bld_base,
args.out[2] = base->undef; /* B, sample mask */
args.out[3] = base->undef; /* A, alpha to mask */
if (format == V_028710_SPI_SHADER_UINT16_ABGR) {
assert(!depth);
args.compr = 1; /* COMPR flag */
if (stencil) {
/* Stencil should be in X[23:16]. */
stencil = ac_to_integer(&ctx->ac, stencil);
- stencil = LLVMBuildShl(ctx->gallivm.builder, stencil,
+ stencil = LLVMBuildShl(ctx->ac.builder, stencil,
LLVMConstInt(ctx->i32, 16, 0), "");
args.out[0] = ac_to_float(&ctx->ac, stencil);
mask |= 0x3;
}
if (samplemask) {
/* SampleMask should be in Y[15:0]. */
args.out[1] = samplemask;
mask |= 0xc;
}
} else {
@@ -3476,21 +3454,21 @@ static void si_export_null(struct lp_build_tgsi_context *bld_base)
*
* The alpha-ref SGPR is returned via its original location.
*/
static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi,
unsigned max_outputs,
LLVMValueRef *addrs)
{
struct si_shader_context *ctx = si_shader_context_from_abi(abi);
struct si_shader *shader = ctx->shader;
struct tgsi_shader_info *info = &shader->selector->info;
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
unsigned i, j, first_vgpr, vgpr;
LLVMValueRef color[8][4] = {};
LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
LLVMValueRef ret;
if (ctx->postponed_kill)
ac_build_kill(&ctx->ac, LLVMBuildLoad(builder, ctx->postponed_kill, ""));
/* Read the output values. */
@@ -3556,26 +3534,24 @@ static void si_llvm_return_fs_outputs(struct ac_shader_abi *abi,
vgpr = first_vgpr + PS_EPILOG_SAMPLEMASK_MIN_LOC;
ret = LLVMBuildInsertValue(builder, ret,
LLVMGetParam(ctx->main_fn,
SI_PARAM_SAMPLE_COVERAGE), vgpr++, "");
ctx->return_value = ret;
}
void si_emit_waitcnt(struct si_shader_context *ctx, unsigned simm16)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef args[1] = {
LLVMConstInt(ctx->i32, simm16, 0)
};
- lp_build_intrinsic(builder, "llvm.amdgcn.s.waitcnt",
+ lp_build_intrinsic(ctx->ac.builder, "llvm.amdgcn.s.waitcnt",
ctx->voidt, args, 1, 0);
}
static void membar_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef src0 = lp_build_emit_fetch(bld_base, emit_data->inst, 0, 0);
@@ -3596,31 +3572,30 @@ static void membar_emit(
if (waitcnt != NOOP_WAITCNT)
si_emit_waitcnt(ctx, waitcnt);
}
static void clock_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef tmp;
- tmp = lp_build_intrinsic(gallivm->builder, "llvm.readcyclecounter",
+ tmp = lp_build_intrinsic(ctx->ac.builder, "llvm.readcyclecounter",
ctx->i64, NULL, 0, 0);
- tmp = LLVMBuildBitCast(gallivm->builder, tmp, ctx->v2i32, "");
+ tmp = LLVMBuildBitCast(ctx->ac.builder, tmp, ctx->v2i32, "");
emit_data->output[0] =
- LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_0, "");
+ LLVMBuildExtractElement(ctx->ac.builder, tmp, ctx->i32_0, "");
emit_data->output[1] =
- LLVMBuildExtractElement(gallivm->builder, tmp, ctx->i32_1, "");
+ LLVMBuildExtractElement(ctx->ac.builder, tmp, ctx->i32_1, "");
}
LLVMTypeRef si_const_array(LLVMTypeRef elem_type, int num_elements)
{
return LLVMPointerType(LLVMArrayType(elem_type, num_elements),
CONST_ADDR_SPACE);
}
static void si_llvm_emit_ddxy(
const struct lp_build_tgsi_action *action,
@@ -3656,21 +3631,21 @@ static void si_llvm_emit_ddxy(
static LLVMValueRef si_llvm_emit_ddxy_interp(
struct lp_build_tgsi_context *bld_base,
LLVMValueRef interp_ij)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef result[4], a;
unsigned i;
for (i = 0; i < 2; i++) {
- a = LLVMBuildExtractElement(gallivm->builder, interp_ij,
+ a = LLVMBuildExtractElement(ctx->ac.builder, interp_ij,
LLVMConstInt(ctx->i32, i, 0), "");
result[i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDX, a);
result[2+i] = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_DDY, a);
}
return lp_build_gather_values(gallivm, result, 4);
}
static void interp_fetch_args(
struct lp_build_tgsi_context *bld_base,
@@ -3721,29 +3696,29 @@ static void interp_fetch_args(
LLVMConstReal(ctx->f32, 0.5),
ctx->ac.f32_0,
ctx->ac.f32_0,
};
sample_position = lp_build_gather_values(gallivm, center, 4);
} else {
sample_position = load_sample_position(ctx, sample_id);
}
- emit_data->args[0] = LLVMBuildExtractElement(gallivm->builder,
+ emit_data->args[0] = LLVMBuildExtractElement(ctx->ac.builder,
sample_position,
ctx->i32_0, "");
- emit_data->args[0] = LLVMBuildFSub(gallivm->builder, emit_data->args[0], halfval, "");
- emit_data->args[1] = LLVMBuildExtractElement(gallivm->builder,
+ emit_data->args[0] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[0], halfval, "");
+ emit_data->args[1] = LLVMBuildExtractElement(ctx->ac.builder,
sample_position,
ctx->i32_1, "");
- emit_data->args[1] = LLVMBuildFSub(gallivm->builder, emit_data->args[1], halfval, "");
+ emit_data->args[1] = LLVMBuildFSub(ctx->ac.builder, emit_data->args[1], halfval, "");
emit_data->arg_count = 2;
}
}
static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct si_shader *shader = ctx->shader;
@@ -3807,115 +3782,112 @@ static void build_interp_intrinsic(const struct lp_build_tgsi_action *action,
* take the I then J parameters, and the DDX/Y for it, and
* calculate the IJ inputs for the interpolator.
* temp1 = ddx * offset/sample.x + I;
* interp_param.I = ddy * offset/sample.y + temp1;
* temp1 = ddx * offset/sample.x + J;
* interp_param.J = ddy * offset/sample.y + temp1;
*/
for (i = 0; i < 2; i++) {
LLVMValueRef ix_ll = LLVMConstInt(ctx->i32, i, 0);
LLVMValueRef iy_ll = LLVMConstInt(ctx->i32, i + 2, 0);
- LLVMValueRef ddx_el = LLVMBuildExtractElement(gallivm->builder,
+ LLVMValueRef ddx_el = LLVMBuildExtractElement(ctx->ac.builder,
ddxy_out, ix_ll, "");
- LLVMValueRef ddy_el = LLVMBuildExtractElement(gallivm->builder,
+ LLVMValueRef ddy_el = LLVMBuildExtractElement(ctx->ac.builder,
ddxy_out, iy_ll, "");
- LLVMValueRef interp_el = LLVMBuildExtractElement(gallivm->builder,
+ LLVMValueRef interp_el = LLVMBuildExtractElement(ctx->ac.builder,
interp_param, ix_ll, "");
LLVMValueRef temp1, temp2;
interp_el = ac_to_float(&ctx->ac, interp_el);
- temp1 = LLVMBuildFMul(gallivm->builder, ddx_el, emit_data->args[0], "");
+ temp1 = LLVMBuildFMul(ctx->ac.builder, ddx_el, emit_data->args[0], "");
- temp1 = LLVMBuildFAdd(gallivm->builder, temp1, interp_el, "");
+ temp1 = LLVMBuildFAdd(ctx->ac.builder, temp1, interp_el, "");
- temp2 = LLVMBuildFMul(gallivm->builder, ddy_el, emit_data->args[1], "");
+ temp2 = LLVMBuildFMul(ctx->ac.builder, ddy_el, emit_data->args[1], "");
- ij_out[i] = LLVMBuildFAdd(gallivm->builder, temp2, temp1, "");
+ ij_out[i] = LLVMBuildFAdd(ctx->ac.builder, temp2, temp1, "");
}
interp_param = lp_build_gather_values(gallivm, ij_out, 2);
}
if (interp_param)
interp_param = ac_to_float(&ctx->ac, interp_param);
for (chan = 0; chan < 4; chan++) {
LLVMValueRef gather = LLVMGetUndef(LLVMVectorType(ctx->f32, input_array_size));
unsigned schan = tgsi_util_get_full_src_register_swizzle(&inst->Src[0], chan);
for (unsigned idx = 0; idx < input_array_size; ++idx) {
LLVMValueRef v, i = NULL, j = NULL;
if (interp_param) {
i = LLVMBuildExtractElement(
- gallivm->builder, interp_param, ctx->i32_0, "");
+ ctx->ac.builder, interp_param, ctx->i32_0, "");
j = LLVMBuildExtractElement(
- gallivm->builder, interp_param, ctx->i32_1, "");
+ ctx->ac.builder, interp_param, ctx->i32_1, "");
}
v = si_build_fs_interp(ctx, input_base + idx, schan,
prim_mask, i, j);
- gather = LLVMBuildInsertElement(gallivm->builder,
+ gather = LLVMBuildInsertElement(ctx->ac.builder,
gather, v, LLVMConstInt(ctx->i32, idx, false), "");
}
emit_data->output[chan] = LLVMBuildExtractElement(
- gallivm->builder, gather, array_idx, "");
+ ctx->ac.builder, gather, array_idx, "");
}
}
static void vote_all_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef tmp = ac_build_vote_all(&ctx->ac, emit_data->args[0]);
emit_data->output[emit_data->chan] =
- LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+ LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, "");
}
static void vote_any_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef tmp = ac_build_vote_any(&ctx->ac, emit_data->args[0]);
emit_data->output[emit_data->chan] =
- LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+ LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, "");
}
static void vote_eq_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef tmp = ac_build_vote_eq(&ctx->ac, emit_data->args[0]);
emit_data->output[emit_data->chan] =
- LLVMBuildSExt(gallivm->builder, tmp, ctx->i32, "");
+ LLVMBuildSExt(ctx->ac.builder, tmp, ctx->i32, "");
}
static void ballot_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef tmp;
tmp = lp_build_emit_fetch(bld_base, emit_data->inst, 0, TGSI_CHAN_X);
tmp = ac_build_ballot(&ctx->ac, tmp);
tmp = LLVMBuildBitCast(builder, tmp, ctx->v2i32, "");
emit_data->output[0] = LLVMBuildExtractElement(builder, tmp, ctx->i32_0, "");
emit_data->output[1] = LLVMBuildExtractElement(builder, tmp, ctx->i32_1, "");
}
@@ -3985,33 +3957,33 @@ static void si_llvm_emit_vertex(
ctx->param_gs2vs_offset);
LLVMValueRef gs_next_vertex;
LLVMValueRef can_emit, kill;
unsigned chan, offset;
int i;
unsigned stream;
stream = si_llvm_get_stream(bld_base, emit_data);
/* Write vertex attribute values to GSVS ring */
- gs_next_vertex = LLVMBuildLoad(gallivm->builder,
+ gs_next_vertex = LLVMBuildLoad(ctx->ac.builder,
ctx->gs_next_vertex[stream],
"");
/* If this thread has already emitted the declared maximum number of
* vertices, skip the write: excessive vertex emissions are not
* supposed to have any effect.
*
* If the shader has no writes to memory, kill it instead. This skips
* further memory loads and may allow LLVM to skip to the end
* altogether.
*/
- can_emit = LLVMBuildICmp(gallivm->builder, LLVMIntULT, gs_next_vertex,
+ can_emit = LLVMBuildICmp(ctx->ac.builder, LLVMIntULT, gs_next_vertex,
LLVMConstInt(ctx->i32,
shader->selector->gs_max_out_vertices, 0), "");
bool use_kill = !info->writes_memory;
if (use_kill) {
kill = lp_build_select(&bld_base->base, can_emit,
LLVMConstReal(ctx->f32, 1.0f),
LLVMConstReal(ctx->f32, -1.0f));
ac_build_kill(&ctx->ac, kill);
@@ -4021,21 +3993,21 @@ static void si_llvm_emit_vertex(
offset = 0;
for (i = 0; i < info->num_outputs; i++) {
LLVMValueRef *out_ptr = ctx->outputs[i];
for (chan = 0; chan < 4; chan++) {
if (!(info->output_usagemask[i] & (1 << chan)) ||
((info->output_streams[i] >> (2 * chan)) & 3) != stream)
continue;
- LLVMValueRef out_val = LLVMBuildLoad(gallivm->builder, out_ptr[chan], "");
+ LLVMValueRef out_val = LLVMBuildLoad(ctx->ac.builder, out_ptr[chan], "");
LLVMValueRef voffset =
LLVMConstInt(ctx->i32, offset *
shader->selector->gs_max_out_vertices, 0);
offset++;
voffset = lp_build_add(uint, voffset, gs_next_vertex);
voffset = lp_build_mul_imm(uint, voffset, 4);
out_val = ac_to_integer(&ctx->ac, out_val);
@@ -4043,21 +4015,21 @@ static void si_llvm_emit_vertex(
ctx->gsvs_ring[stream],
out_val, 1,
voffset, soffset, 0,
1, 1, true, true);
}
}
gs_next_vertex = lp_build_add(uint, gs_next_vertex,
ctx->i32_1);
- LLVMBuildStore(gallivm->builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
+ LLVMBuildStore(ctx->ac.builder, gs_next_vertex, ctx->gs_next_vertex[stream]);
/* Signal vertex emission */
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_EMIT | AC_SENDMSG_GS | (stream << 8),
si_get_gs_wave_id(ctx));
if (!use_kill)
lp_build_endif(&if_state);
}
/* Cut one primitive from the geometry shader */
static void si_llvm_emit_primitive(
@@ -4072,33 +4044,32 @@ static void si_llvm_emit_primitive(
stream = si_llvm_get_stream(bld_base, emit_data);
ac_build_sendmsg(&ctx->ac, AC_SENDMSG_GS_OP_CUT | AC_SENDMSG_GS | (stream << 8),
si_get_gs_wave_id(ctx));
}
static void si_llvm_emit_barrier(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
/* SI only (thanks to a hw bug workaround):
* The real barrier instruction isn’t needed, because an entire patch
* always fits into a single wave.
*/
if (ctx->screen->b.chip_class == SI &&
ctx->type == PIPE_SHADER_TESS_CTRL) {
si_emit_waitcnt(ctx, LGKM_CNT & VM_CNT);
return;
}
- lp_build_intrinsic(gallivm->builder,
+ lp_build_intrinsic(ctx->ac.builder,
"llvm.amdgcn.s.barrier",
ctx->voidt, NULL, 0, LP_FUNC_ATTR_CONVERGENT);
}
static const struct lp_build_tgsi_action interp_action = {
.fetch_args = interp_fetch_args,
.emit = build_interp_intrinsic,
};
static void si_create_function(struct si_shader_context *ctx,
@@ -4180,24 +4151,22 @@ static void declare_streamout_params(struct si_shader_context *ctx,
for (i = 0; i < 4; i++) {
if (!so->stride[i])
continue;
ctx->param_streamout_offset[i] = add_arg(fninfo, ARG_SGPR, ctx->ac.i32);
}
}
static void declare_lds_as_pointer(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
-
unsigned lds_size = ctx->screen->b.chip_class >= CIK ? 65536 : 32768;
- ctx->lds = LLVMBuildIntToPtr(gallivm->builder, ctx->i32_0,
+ ctx->lds = LLVMBuildIntToPtr(ctx->ac.builder, ctx->i32_0,
LLVMPointerType(LLVMArrayType(ctx->i32, lds_size / 4), LOCAL_ADDR_SPACE),
"lds");
}
static unsigned si_get_max_workgroup_size(const struct si_shader *shader)
{
switch (shader->selector->type) {
case PIPE_SHADER_TESS_CTRL:
/* Return this so that LLVM doesn't remove s_barrier
* instructions on chips where we use s_barrier. */
@@ -4656,22 +4625,21 @@ static void create_function(struct si_shader_context *ctx)
ctx->type == PIPE_SHADER_GEOMETRY)))
declare_lds_as_pointer(ctx);
}
/**
* Load ESGS and GSVS ring buffer resource descriptors and save the variables
* for later use.
*/
static void preload_ring_buffers(struct si_shader_context *ctx)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef buf_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_rw_buffers);
if (ctx->screen->b.chip_class <= VI &&
(ctx->shader->key.as_es || ctx->type == PIPE_SHADER_GEOMETRY)) {
unsigned ring =
ctx->type == PIPE_SHADER_GEOMETRY ? SI_GS_RING_ESGS
: SI_ES_RING_ESGS;
LLVMValueRef offset = LLVMConstInt(ctx->i32, ring, 0);
@@ -4754,22 +4722,21 @@ static void preload_ring_buffers(struct si_shader_context *ctx)
ctx->gsvs_ring[stream] = ring;
}
}
}
static void si_llvm_emit_polygon_stipple(struct si_shader_context *ctx,
LLVMValueRef param_rw_buffers,
unsigned param_pos_fixed_pt)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef slot, desc, offset, row, bit, address[2];
/* Use the fixed-point gl_FragCoord input.
* Since the stipple pattern is 32x32 and it repeats, just get 5 bits
* per coordinate to get the repeating effect.
*/
address[0] = unpack_param(ctx, param_pos_fixed_pt, 0, 5);
address[1] = unpack_param(ctx, param_pos_fixed_pt, 16, 5);
/* Load the buffer descriptor. */
@@ -5272,23 +5239,23 @@ static int si_compile_llvm(struct si_screen *sscreen,
fprintf(stderr, "radeonsi: The shader can't have rodata.");
return -EINVAL;
}
return r;
}
static void si_llvm_build_ret(struct si_shader_context *ctx, LLVMValueRef ret)
{
if (LLVMGetTypeKind(LLVMTypeOf(ret)) == LLVMVoidTypeKind)
- LLVMBuildRetVoid(ctx->gallivm.builder);
+ LLVMBuildRetVoid(ctx->ac.builder);
else
- LLVMBuildRet(ctx->gallivm.builder, ret);
+ LLVMBuildRet(ctx->ac.builder, ret);
}
/* Generate code for the hardware VS shader stage to go with a geometry shader */
struct si_shader *
si_generate_gs_copy_shader(struct si_screen *sscreen,
LLVMTargetMachineRef tm,
struct si_shader_selector *gs_selector,
struct pipe_debug_callback *debug)
{
struct si_shader_context ctx;
@@ -5313,21 +5280,21 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
}
shader->selector = gs_selector;
shader->is_gs_copy_shader = true;
si_init_shader_ctx(&ctx, sscreen, tm);
ctx.shader = shader;
ctx.type = PIPE_SHADER_VERTEX;
- builder = gallivm->builder;
+ builder = ctx.ac.builder;
create_function(&ctx);
preload_ring_buffers(&ctx);
LLVMValueRef voffset =
lp_build_mul_imm(uint, ctx.abi.vertex_id, 4);
/* Fetch the vertex stream ID.*/
LLVMValueRef stream_id;
@@ -5398,21 +5365,21 @@ si_generate_gs_copy_shader(struct si_screen *sscreen,
}
if (stream == 0)
si_llvm_export_vs(bld_base, outputs, gsinfo->num_outputs);
LLVMBuildBr(builder, end_bb);
}
LLVMPositionBuilderAtEnd(builder, end_bb);
- LLVMBuildRetVoid(gallivm->builder);
+ LLVMBuildRetVoid(ctx.ac.builder);
ctx.type = PIPE_SHADER_GEOMETRY; /* override for shader dumping */
si_llvm_optimize_module(&ctx);
r = si_compile_llvm(sscreen, &ctx.shader->binary,
&ctx.shader->config, ctx.tm,
ctx.gallivm.module,
debug, PIPE_SHADER_GEOMETRY,
"GS Copy Shader");
if (!r) {
@@ -5610,33 +5577,33 @@ static void si_count_scratch_private_memory(struct si_shader_context *ctx)
unsigned dw_size = align(ac_get_type_size(type) / 4, alignment);
ctx->shader->config.private_mem_vgprs += dw_size;
}
bb = LLVMGetNextBasicBlock(bb);
}
}
static void si_init_exec_full_mask(struct si_shader_context *ctx)
{
LLVMValueRef full_mask = LLVMConstInt(ctx->i64, ~0ull, 0);
- lp_build_intrinsic(ctx->gallivm.builder,
+ lp_build_intrinsic(ctx->ac.builder,
"llvm.amdgcn.init.exec", ctx->voidt,
&full_mask, 1, LP_FUNC_ATTR_CONVERGENT);
}
static void si_init_exec_from_input(struct si_shader_context *ctx,
unsigned param, unsigned bitoffset)
{
LLVMValueRef args[] = {
LLVMGetParam(ctx->main_fn, param),
LLVMConstInt(ctx->i32, bitoffset, 0),
};
- lp_build_intrinsic(ctx->gallivm.builder,
+ lp_build_intrinsic(ctx->ac.builder,
"llvm.amdgcn.init.exec.from.input",
ctx->voidt, args, 2, LP_FUNC_ATTR_CONVERGENT);
}
static bool si_vs_needs_prolog(const struct si_shader_selector *sel,
const struct si_vs_prolog_bits *key)
{
/* VGPR initialization fixup for Vega10 and Raven is always done in the
* VS prolog. */
return sel->vs_needs_prolog || key->ls_vgpr_fix;
@@ -5968,23 +5935,22 @@ static void si_get_ps_epilog_key(struct si_shader *shader,
}
/**
* Build the GS prolog function. Rotate the input vertices for triangle strips
* with adjacency.
*/
static void si_build_gs_prolog_function(struct si_shader_context *ctx,
union si_shader_part_key *key)
{
unsigned num_sgprs, num_vgprs;
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_function_info fninfo;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMTypeRef returns[48];
LLVMValueRef func, ret;
si_init_function_info(&fninfo);
if (ctx->screen->b.chip_class >= GFX9) {
num_sgprs = 8 + GFX9_GS_NUM_USER_SGPR;
num_vgprs = 5; /* ES inputs are not needed by GS */
} else {
num_sgprs = GFX6_GS_NUM_USER_SGPR + 2;
@@ -6094,21 +6060,21 @@ static void si_build_gs_prolog_function(struct si_shader_context *ctx,
* Given a list of shader part functions, build a wrapper function that
* runs them in sequence to form a monolithic shader.
*/
static void si_build_wrapper_function(struct si_shader_context *ctx,
LLVMValueRef *parts,
unsigned num_parts,
unsigned main_part,
unsigned next_shader_first_part)
{
struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
/* PS epilog has one arg per color component; gfx9 merged shader
* prologs need to forward 32 user SGPRs.
*/
struct si_function_info fninfo;
LLVMValueRef initial[64], out[64];
LLVMTypeRef function_type;
unsigned num_first_params;
unsigned num_out, initial_num_out;
MAYBE_UNUSED unsigned num_out_sgpr; /* used in debug checks */
MAYBE_UNUSED unsigned initial_num_out_sgpr; /* used in debug checks */
@@ -6727,22 +6693,22 @@ out:
static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
{
struct gallivm_state *gallivm = &ctx->gallivm;
LLVMValueRef ptr[2], list;
/* Get the pointer to rw buffers. */
ptr[0] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS);
ptr[1] = LLVMGetParam(ctx->main_fn, SI_SGPR_RW_BUFFERS_HI);
list = lp_build_gather_values(gallivm, ptr, 2);
- list = LLVMBuildBitCast(gallivm->builder, list, ctx->i64, "");
- list = LLVMBuildIntToPtr(gallivm->builder, list,
+ list = LLVMBuildBitCast(ctx->ac.builder, list, ctx->i64, "");
+ list = LLVMBuildIntToPtr(ctx->ac.builder, list,
si_const_array(ctx->v4i32, SI_NUM_RW_BUFFERS), "");
return list;
}
/**
* Build the vertex shader prolog function.
*
* The inputs are the same as VS (a lot of SGPRs and 4 VGPR system values).
* All inputs are returned unmodified. The vertex load indices are
* stored after them, which will be used by the API VS for fetching inputs.
@@ -6752,21 +6718,20 @@ static LLVMValueRef si_prolog_get_rw_buffers(struct si_shader_context *ctx)
* input_v1,
* input_v2,
* input_v3,
* (VertexID + BaseVertex),
* (InstanceID + StartInstance),
* (InstanceID / 2 + StartInstance)
*/
static void si_build_vs_prolog_function(struct si_shader_context *ctx,
union si_shader_part_key *key)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_function_info fninfo;
LLVMTypeRef *returns;
LLVMValueRef ret, func;
int num_returns, i;
unsigned first_vs_vgpr = key->vs_prolog.num_merged_next_stage_vgprs;
unsigned num_input_vgprs = key->vs_prolog.num_merged_next_stage_vgprs + 4;
LLVMValueRef input_vgprs[9];
unsigned num_all_input_regs = key->vs_prolog.num_input_sgprs +
num_input_vgprs;
unsigned user_sgpr_base = key->vs_prolog.num_merged_next_stage_vgprs ? 8 : 0;
@@ -6803,48 +6768,48 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
si_init_exec_from_input(ctx, 3, 0);
if (key->vs_prolog.as_ls &&
(ctx->screen->b.family == CHIP_VEGA10 ||
ctx->screen->b.family == CHIP_RAVEN)) {
/* If there are no HS threads, SPI loads the LS VGPRs
* starting at VGPR 0. Shift them back to where they
* belong.
*/
LLVMValueRef has_hs_threads =
- LLVMBuildICmp(gallivm->builder, LLVMIntNE,
+ LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
unpack_param(ctx, 3, 8, 8),
ctx->i32_0, "");
for (i = 4; i > 0; --i) {
input_vgprs[i + 1] =
- LLVMBuildSelect(gallivm->builder, has_hs_threads,
+ LLVMBuildSelect(ctx->ac.builder, has_hs_threads,
input_vgprs[i + 1],
input_vgprs[i - 1], "");
}
}
}
ctx->abi.vertex_id = input_vgprs[first_vs_vgpr];
ctx->abi.instance_id = input_vgprs[first_vs_vgpr + (key->vs_prolog.as_ls ? 2 : 1)];
/* Copy inputs to outputs. This should be no-op, as the registers match,
* but it will prevent the compiler from overwriting them unintentionally.
*/
ret = ctx->return_value;
for (i = 0; i < key->vs_prolog.num_input_sgprs; i++) {
LLVMValueRef p = LLVMGetParam(func, i);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, p, i, "");
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p, i, "");
}
for (i = 0; i < num_input_vgprs; i++) {
LLVMValueRef p = input_vgprs[i];
p = ac_to_float(&ctx->ac, p);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, p,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p,
key->vs_prolog.num_input_sgprs + i, "");
}
/* Compute vertex load indices from instance divisors. */
LLVMValueRef instance_divisor_constbuf = NULL;
if (key->vs_prolog.states.instance_divisor_is_fetched) {
LLVMValueRef list = si_prolog_get_rw_buffers(ctx);
LLVMValueRef buf_index =
LLVMConstInt(ctx->i32, SI_VS_CONST_INSTANCE_DIVISORS, 0);
@@ -6868,28 +6833,28 @@ static void si_build_vs_prolog_function(struct si_shader_context *ctx,
divisor = ac_to_integer(&ctx->ac, divisor);
}
/* InstanceID / Divisor + StartInstance */
index = get_instance_index_for_fetch(ctx,
user_sgpr_base +
SI_SGPR_START_INSTANCE,
divisor);
} else {
/* VertexID + BaseVertex */
- index = LLVMBuildAdd(gallivm->builder,
+ index = LLVMBuildAdd(ctx->ac.builder,
ctx->abi.vertex_id,
LLVMGetParam(func, user_sgpr_base +
SI_SGPR_BASE_VERTEX), "");
}
index = ac_to_float(&ctx->ac, index);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, index,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, index,
fninfo.num_params + i, "");
}
si_llvm_build_ret(ctx, ret);
}
static bool si_get_vs_prolog(struct si_screen *sscreen,
LLVMTargetMachineRef tm,
struct si_shader *shader,
struct pipe_debug_callback *debug,
@@ -6926,21 +6891,20 @@ static bool si_shader_select_vs_parts(struct si_screen *sscreen,
&shader->key.part.vs.prolog);
}
/**
* Compile the TCS epilog function. This writes tesselation factors to memory
* based on the output primitive type of the tesselator (determined by TES).
*/
static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
union si_shader_part_key *key)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
struct si_function_info fninfo;
LLVMValueRef func;
si_init_function_info(&fninfo);
if (ctx->screen->b.chip_class >= GFX9) {
add_arg(&fninfo, ARG_SGPR, ctx->i64);
ctx->param_tcs_offchip_offset = add_arg(&fninfo, ARG_SGPR, ctx->i32);
add_arg(&fninfo, ARG_SGPR, ctx->i32); /* wave info */
@@ -6996,21 +6960,21 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx,
LLVMValueRef invoc0_tess_factors[6];
for (unsigned i = 0; i < 6; i++)
invoc0_tess_factors[i] = LLVMGetParam(func, tess_factors_idx + 3 + i);
si_write_tess_factors(bld_base,
LLVMGetParam(func, tess_factors_idx),
LLVMGetParam(func, tess_factors_idx + 1),
LLVMGetParam(func, tess_factors_idx + 2),
invoc0_tess_factors, invoc0_tess_factors + 4);
- LLVMBuildRetVoid(gallivm->builder);
+ LLVMBuildRetVoid(ctx->ac.builder);
}
/**
* Select and compile (or reuse) TCS parts (epilog).
*/
static bool si_shader_select_tcs_parts(struct si_screen *sscreen,
LLVMTargetMachineRef tm,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
@@ -7112,21 +7076,21 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
si_create_function(ctx, "ps_prolog", fninfo.types, num_returns,
&fninfo, 0);
func = ctx->main_fn;
/* Copy inputs to outputs. This should be no-op, as the registers match,
* but it will prevent the compiler from overwriting them unintentionally.
*/
ret = ctx->return_value;
for (i = 0; i < fninfo.num_params; i++) {
LLVMValueRef p = LLVMGetParam(func, i);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, p, i, "");
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, p, i, "");
}
/* Polygon stippling. */
if (key->ps_prolog.states.poly_stipple) {
/* POS_FIXED_PT is always last. */
unsigned pos = key->ps_prolog.num_input_sgprs +
key->ps_prolog.num_input_vgprs - 1;
LLVMValueRef list = si_prolog_get_rw_buffers(ctx);
si_llvm_emit_polygon_stipple(ctx, list, pos);
@@ -7137,122 +7101,122 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
unsigned i, base = key->ps_prolog.num_input_sgprs;
LLVMValueRef center[2], centroid[2], tmp, bc_optimize;
/* The shader should do: if (PRIM_MASK[31]) CENTROID = CENTER;
* The hw doesn't compute CENTROID if the whole wave only
* contains fully-covered quads.
*
* PRIM_MASK is after user SGPRs.
*/
bc_optimize = LLVMGetParam(func, SI_PS_NUM_USER_SGPR);
- bc_optimize = LLVMBuildLShr(gallivm->builder, bc_optimize,
+ bc_optimize = LLVMBuildLShr(ctx->ac.builder, bc_optimize,
LLVMConstInt(ctx->i32, 31, 0), "");
- bc_optimize = LLVMBuildTrunc(gallivm->builder, bc_optimize,
+ bc_optimize = LLVMBuildTrunc(ctx->ac.builder, bc_optimize,
ctx->i1, "");
if (key->ps_prolog.states.bc_optimize_for_persp) {
/* Read PERSP_CENTER. */
for (i = 0; i < 2; i++)
center[i] = LLVMGetParam(func, base + 2 + i);
/* Read PERSP_CENTROID. */
for (i = 0; i < 2; i++)
centroid[i] = LLVMGetParam(func, base + 4 + i);
/* Select PERSP_CENTROID. */
for (i = 0; i < 2; i++) {
- tmp = LLVMBuildSelect(gallivm->builder, bc_optimize,
+ tmp = LLVMBuildSelect(ctx->ac.builder, bc_optimize,
center[i], centroid[i], "");
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
tmp, base + 4 + i, "");
}
}
if (key->ps_prolog.states.bc_optimize_for_linear) {
/* Read LINEAR_CENTER. */
for (i = 0; i < 2; i++)
center[i] = LLVMGetParam(func, base + 8 + i);
/* Read LINEAR_CENTROID. */
for (i = 0; i < 2; i++)
centroid[i] = LLVMGetParam(func, base + 10 + i);
/* Select LINEAR_CENTROID. */
for (i = 0; i < 2; i++) {
- tmp = LLVMBuildSelect(gallivm->builder, bc_optimize,
+ tmp = LLVMBuildSelect(ctx->ac.builder, bc_optimize,
center[i], centroid[i], "");
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
tmp, base + 10 + i, "");
}
}
}
/* Force per-sample interpolation. */
if (key->ps_prolog.states.force_persp_sample_interp) {
unsigned i, base = key->ps_prolog.num_input_sgprs;
LLVMValueRef persp_sample[2];
/* Read PERSP_SAMPLE. */
for (i = 0; i < 2; i++)
persp_sample[i] = LLVMGetParam(func, base + i);
/* Overwrite PERSP_CENTER. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
persp_sample[i], base + 2 + i, "");
/* Overwrite PERSP_CENTROID. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
persp_sample[i], base + 4 + i, "");
}
if (key->ps_prolog.states.force_linear_sample_interp) {
unsigned i, base = key->ps_prolog.num_input_sgprs;
LLVMValueRef linear_sample[2];
/* Read LINEAR_SAMPLE. */
for (i = 0; i < 2; i++)
linear_sample[i] = LLVMGetParam(func, base + 6 + i);
/* Overwrite LINEAR_CENTER. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
linear_sample[i], base + 8 + i, "");
/* Overwrite LINEAR_CENTROID. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
linear_sample[i], base + 10 + i, "");
}
/* Force center interpolation. */
if (key->ps_prolog.states.force_persp_center_interp) {
unsigned i, base = key->ps_prolog.num_input_sgprs;
LLVMValueRef persp_center[2];
/* Read PERSP_CENTER. */
for (i = 0; i < 2; i++)
persp_center[i] = LLVMGetParam(func, base + 2 + i);
/* Overwrite PERSP_SAMPLE. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
persp_center[i], base + i, "");
/* Overwrite PERSP_CENTROID. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
persp_center[i], base + 4 + i, "");
}
if (key->ps_prolog.states.force_linear_center_interp) {
unsigned i, base = key->ps_prolog.num_input_sgprs;
LLVMValueRef linear_center[2];
/* Read LINEAR_CENTER. */
for (i = 0; i < 2; i++)
linear_center[i] = LLVMGetParam(func, base + 8 + i);
/* Overwrite LINEAR_SAMPLE. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
linear_center[i], base + 6 + i, "");
/* Overwrite LINEAR_CENTROID. */
for (i = 0; i < 2; i++)
- ret = LLVMBuildInsertValue(gallivm->builder, ret,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret,
linear_center[i], base + 10 + i, "");
}
/* Interpolate colors. */
unsigned color_out_idx = 0;
for (i = 0; i < 2; i++) {
unsigned writemask = (key->ps_prolog.colors_read >> (i * 4)) & 0xf;
unsigned face_vgpr = key->ps_prolog.num_input_sgprs +
key->ps_prolog.face_vgpr_index;
LLVMValueRef interp[2], color[4];
@@ -7260,23 +7224,23 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
if (!writemask)
continue;
/* If the interpolation qualifier is not CONSTANT (-1). */
if (key->ps_prolog.color_interp_vgpr_index[i] != -1) {
unsigned interp_vgpr = key->ps_prolog.num_input_sgprs +
key->ps_prolog.color_interp_vgpr_index[i];
/* Get the (i,j) updated by bc_optimize handling. */
- interp[0] = LLVMBuildExtractValue(gallivm->builder, ret,
+ interp[0] = LLVMBuildExtractValue(ctx->ac.builder, ret,
interp_vgpr, "");
- interp[1] = LLVMBuildExtractValue(gallivm->builder, ret,
+ interp[1] = LLVMBuildExtractValue(ctx->ac.builder, ret,
interp_vgpr + 1, "");
interp_ij = lp_build_gather_values(gallivm, interp, 2);
}
/* Use the absolute location of the input. */
prim_mask = LLVMGetParam(func, SI_PS_NUM_USER_SGPR);
if (key->ps_prolog.states.color_two_side) {
face = LLVMGetParam(func, face_vgpr);
face = ac_to_integer(&ctx->ac, face);
@@ -7284,21 +7248,21 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
interp_fs_input(ctx,
key->ps_prolog.color_attr_index[i],
TGSI_SEMANTIC_COLOR, i,
key->ps_prolog.num_interp_inputs,
key->ps_prolog.colors_read, interp_ij,
prim_mask, face, color);
while (writemask) {
unsigned chan = u_bit_scan(&writemask);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, color[chan],
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, color[chan],
fninfo.num_params + color_out_idx++, "");
}
}
/* Section 15.2.2 (Shader Inputs) of the OpenGL 4.5 (Core Profile) spec
* says:
*
* "When per-sample shading is active due to the use of a fragment
* input qualified by sample or due to the use of the gl_SampleID
* or gl_SamplePosition variables, only the bit for the current
@@ -7325,49 +7289,48 @@ static void si_build_ps_prolog_function(struct si_shader_context *ctx,
assert(key->ps_prolog.states.samplemask_log_ps_iter < ARRAY_SIZE(ps_iter_masks));
uint32_t ps_iter_mask = ps_iter_masks[key->ps_prolog.states.samplemask_log_ps_iter];
unsigned ancillary_vgpr = key->ps_prolog.num_input_sgprs +
key->ps_prolog.ancillary_vgpr_index;
LLVMValueRef sampleid = unpack_param(ctx, ancillary_vgpr, 8, 4);
LLVMValueRef samplemask = LLVMGetParam(func, ancillary_vgpr + 1);
samplemask = ac_to_integer(&ctx->ac, samplemask);
samplemask = LLVMBuildAnd(
- gallivm->builder,
+ ctx->ac.builder,
samplemask,
- LLVMBuildShl(gallivm->builder,
+ LLVMBuildShl(ctx->ac.builder,
LLVMConstInt(ctx->i32, ps_iter_mask, false),
sampleid, ""),
"");
samplemask = ac_to_float(&ctx->ac, samplemask);
- ret = LLVMBuildInsertValue(gallivm->builder, ret, samplemask,
+ ret = LLVMBuildInsertValue(ctx->ac.builder, ret, samplemask,
ancillary_vgpr + 1, "");
}
/* Tell LLVM to insert WQM instruction sequence when needed. */
if (key->ps_prolog.wqm) {
LLVMAddTargetDependentFunctionAttr(func,
"amdgpu-ps-wqm-outputs", "");
}
si_llvm_build_ret(ctx, ret);
}
/**
* Build the pixel shader epilog function. This handles everything that must be
* emulated for pixel shader exports. (alpha-test, format conversions, etc)
*/
static void si_build_ps_epilog_function(struct si_shader_context *ctx,
union si_shader_part_key *key)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
struct lp_build_tgsi_context *bld_base = &ctx->bld_base;
struct si_function_info fninfo;
LLVMValueRef depth = NULL, stencil = NULL, samplemask = NULL;
int i;
struct si_ps_exports exp = {};
si_init_function_info(&fninfo);
/* Declare input SGPRs. */
ctx->param_rw_buffers = add_arg(&fninfo, ARG_SGPR, ctx->i64);
@@ -7443,21 +7406,21 @@ static void si_build_ps_epilog_function(struct si_shader_context *ctx,
if (depth || stencil || samplemask)
si_export_mrt_z(bld_base, depth, stencil, samplemask, &exp);
else if (last_color_export == -1)
si_export_null(bld_base);
if (exp.num)
si_emit_ps_exports(ctx, &exp);
/* Compile. */
- LLVMBuildRetVoid(gallivm->builder);
+ LLVMBuildRetVoid(ctx->ac.builder);
}
/**
* Select and compile (or reuse) pixel shader parts (prolog & epilog).
*/
static bool si_shader_select_ps_parts(struct si_screen *sscreen,
LLVMTargetMachineRef tm,
struct si_shader *shader,
struct pipe_debug_callback *debug)
{
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c b/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
index 1048e3b..443f0c9 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_alu.c
@@ -25,22 +25,23 @@
#include "gallivm/lp_bld_const.h"
#include "gallivm/lp_bld_intr.h"
#include "gallivm/lp_bld_gather.h"
#include "tgsi/tgsi_parse.h"
#include "amd/common/ac_llvm_build.h"
static void kill_if_fetch_args(struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
const struct tgsi_full_instruction *inst = emit_data->inst;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
struct gallivm_state *gallivm = bld_base->base.gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
unsigned i;
LLVMValueRef conds[TGSI_NUM_CHANNELS];
for (i = 0; i < TGSI_NUM_CHANNELS; i++) {
LLVMValueRef value = lp_build_emit_fetch(bld_base, inst, 0, i);
conds[i] = LLVMBuildFCmp(builder, LLVMRealOLT, value,
bld_base->base.zero, "");
}
/* Or the conditions together */
@@ -53,21 +54,21 @@ static void kill_if_fetch_args(struct lp_build_tgsi_context *bld_base,
emit_data->args[0] = LLVMBuildSelect(builder, conds[0],
lp_build_const_float(gallivm, -1.0f),
bld_base->base.zero, "");
}
static void kil_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
if (ctx->postponed_kill) {
if (emit_data->inst->Instruction.Opcode == TGSI_OPCODE_KILL_IF) {
LLVMValueRef val;
/* Take the minimum kill value. This is the same as OR
* between 2 kill values. If the value is negative,
* the pixel will be killed.
*/
val = LLVMBuildLoad(builder, ctx->postponed_kill, "");
@@ -86,21 +87,21 @@ static void kil_emit(const struct lp_build_tgsi_action *action,
ac_build_kill(&ctx->ac, emit_data->args[0]);
else
ac_build_kill(&ctx->ac, NULL);
}
static void emit_icmp(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
unsigned pred;
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMContextRef context = bld_base->base.gallivm->context;
switch (emit_data->inst->Instruction.Opcode) {
case TGSI_OPCODE_USEQ:
case TGSI_OPCODE_U64SEQ: pred = LLVMIntEQ; break;
case TGSI_OPCODE_USNE:
case TGSI_OPCODE_U64SNE: pred = LLVMIntNE; break;
case TGSI_OPCODE_USGE:
case TGSI_OPCODE_U64SGE: pred = LLVMIntUGE; break;
case TGSI_OPCODE_USLT:
@@ -108,268 +109,264 @@ static void emit_icmp(const struct lp_build_tgsi_action *action,
case TGSI_OPCODE_ISGE:
case TGSI_OPCODE_I64SGE: pred = LLVMIntSGE; break;
case TGSI_OPCODE_ISLT:
case TGSI_OPCODE_I64SLT: pred = LLVMIntSLT; break;
default:
assert(!"unknown instruction");
pred = 0;
break;
}
- LLVMValueRef v = LLVMBuildICmp(builder, pred,
+ LLVMValueRef v = LLVMBuildICmp(ctx->ac.builder, pred,
emit_data->args[0], emit_data->args[1],"");
- v = LLVMBuildSExtOrBitCast(builder, v,
+ v = LLVMBuildSExtOrBitCast(ctx->ac.builder, v,
LLVMInt32TypeInContext(context), "");
emit_data->output[emit_data->chan] = v;
}
static void emit_ucmp(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
-
LLVMValueRef arg0 = ac_to_integer(&ctx->ac, emit_data->args[0]);
- LLVMValueRef v = LLVMBuildICmp(builder, LLVMIntNE, arg0,
+ LLVMValueRef v = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE, arg0,
ctx->i32_0, "");
emit_data->output[emit_data->chan] =
- LLVMBuildSelect(builder, v, emit_data->args[1], emit_data->args[2], "");
+ LLVMBuildSelect(ctx->ac.builder, v, emit_data->args[1], emit_data->args[2], "");
}
static void emit_cmp(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef cond, *args = emit_data->args;
- cond = LLVMBuildFCmp(builder, LLVMRealOLT, args[0],
+ cond = LLVMBuildFCmp(ctx->ac.builder, LLVMRealOLT, args[0],
bld_base->base.zero, "");
emit_data->output[emit_data->chan] =
- LLVMBuildSelect(builder, cond, args[1], args[2], "");
+ LLVMBuildSelect(ctx->ac.builder, cond, args[1], args[2], "");
}
static void emit_set_cond(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMRealPredicate pred;
LLVMValueRef cond;
/* Use ordered for everything but NE (which is usual for
* float comparisons)
*/
switch (emit_data->inst->Instruction.Opcode) {
case TGSI_OPCODE_SGE: pred = LLVMRealOGE; break;
case TGSI_OPCODE_SEQ: pred = LLVMRealOEQ; break;
case TGSI_OPCODE_SLE: pred = LLVMRealOLE; break;
case TGSI_OPCODE_SLT: pred = LLVMRealOLT; break;
case TGSI_OPCODE_SNE: pred = LLVMRealUNE; break;
case TGSI_OPCODE_SGT: pred = LLVMRealOGT; break;
default: assert(!"unknown instruction"); pred = 0; break;
}
- cond = LLVMBuildFCmp(builder,
+ cond = LLVMBuildFCmp(ctx->ac.builder,
pred, emit_data->args[0], emit_data->args[1], "");
- emit_data->output[emit_data->chan] = LLVMBuildSelect(builder,
+ emit_data->output[emit_data->chan] = LLVMBuildSelect(ctx->ac.builder,
cond, bld_base->base.one, bld_base->base.zero, "");
}
static void emit_fcmp(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMContextRef context = bld_base->base.gallivm->context;
LLVMRealPredicate pred;
/* Use ordered for everything but NE (which is usual for
* float comparisons)
*/
switch (emit_data->inst->Instruction.Opcode) {
case TGSI_OPCODE_FSEQ: pred = LLVMRealOEQ; break;
case TGSI_OPCODE_FSGE: pred = LLVMRealOGE; break;
case TGSI_OPCODE_FSLT: pred = LLVMRealOLT; break;
case TGSI_OPCODE_FSNE: pred = LLVMRealUNE; break;
default: assert(!"unknown instruction"); pred = 0; break;
}
- LLVMValueRef v = LLVMBuildFCmp(builder, pred,
+ LLVMValueRef v = LLVMBuildFCmp(ctx->ac.builder, pred,
emit_data->args[0], emit_data->args[1],"");
- v = LLVMBuildSExtOrBitCast(builder, v,
+ v = LLVMBuildSExtOrBitCast(ctx->ac.builder, v,
LLVMInt32TypeInContext(context), "");
emit_data->output[emit_data->chan] = v;
}
static void emit_dcmp(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMContextRef context = bld_base->base.gallivm->context;
LLVMRealPredicate pred;
/* Use ordered for everything but NE (which is usual for
* float comparisons)
*/
switch (emit_data->inst->Instruction.Opcode) {
case TGSI_OPCODE_DSEQ: pred = LLVMRealOEQ; break;
case TGSI_OPCODE_DSGE: pred = LLVMRealOGE; break;
case TGSI_OPCODE_DSLT: pred = LLVMRealOLT; break;
case TGSI_OPCODE_DSNE: pred = LLVMRealUNE; break;
default: assert(!"unknown instruction"); pred = 0; break;
}
- LLVMValueRef v = LLVMBuildFCmp(builder, pred,
+ LLVMValueRef v = LLVMBuildFCmp(ctx->ac.builder, pred,
emit_data->args[0], emit_data->args[1],"");
- v = LLVMBuildSExtOrBitCast(builder, v,
+ v = LLVMBuildSExtOrBitCast(ctx->ac.builder, v,
LLVMInt32TypeInContext(context), "");
emit_data->output[emit_data->chan] = v;
}
static void emit_not(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
LLVMValueRef v = ac_to_integer(&ctx->ac, emit_data->args[0]);
- emit_data->output[emit_data->chan] = LLVMBuildNot(builder, v, "");
+ emit_data->output[emit_data->chan] = LLVMBuildNot(ctx->ac.builder, v, "");
}
static void emit_arl(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
LLVMValueRef floor_index = lp_build_emit_llvm_unary(bld_base, TGSI_OPCODE_FLR, emit_data->args[0]);
- emit_data->output[emit_data->chan] = LLVMBuildFPToSI(builder,
+ emit_data->output[emit_data->chan] = LLVMBuildFPToSI(ctx->ac.builder,
floor_index, ctx->i32, "");
}
static void emit_and(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildAnd(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildAnd(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_or(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildOr(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildOr(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_uadd(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildAdd(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildAdd(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_udiv(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildUDiv(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildUDiv(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_idiv(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildSDiv(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildSDiv(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_mod(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildSRem(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildSRem(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_umod(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildURem(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildURem(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_shl(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildShl(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildShl(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_ushr(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildLShr(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildLShr(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_ishr(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildAShr(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildAShr(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_xor(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildXor(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildXor(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
}
static void emit_ssg(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef cmp, val;
if (emit_data->inst->Instruction.Opcode == TGSI_OPCODE_I64SSG) {
cmp = LLVMBuildICmp(builder, LLVMIntSGT, emit_data->args[0], bld_base->int64_bld.zero, "");
val = LLVMBuildSelect(builder, cmp, bld_base->int64_bld.one, emit_data->args[0], "");
cmp = LLVMBuildICmp(builder, LLVMIntSGE, val, bld_base->int64_bld.zero, "");
val = LLVMBuildSelect(builder, cmp, val, LLVMConstInt(ctx->i64, -1, true), "");
} else if (emit_data->inst->Instruction.Opcode == TGSI_OPCODE_ISSG) {
cmp = LLVMBuildICmp(builder, LLVMIntSGT, emit_data->args[0], ctx->i32_0, "");
@@ -388,112 +385,112 @@ static void emit_ssg(const struct lp_build_tgsi_action *action,
val = LLVMBuildSelect(builder, cmp, val, LLVMConstReal(ctx->f32, -1), "");
}
emit_data->output[emit_data->chan] = val;
}
static void emit_ineg(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildNeg(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildNeg(ctx->ac.builder,
emit_data->args[0], "");
}
static void emit_dneg(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildFNeg(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildFNeg(ctx->ac.builder,
emit_data->args[0], "");
}
static void emit_frac(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
char *intr;
if (emit_data->info->opcode == TGSI_OPCODE_FRC)
intr = "llvm.floor.f32";
else if (emit_data->info->opcode == TGSI_OPCODE_DFRAC)
intr = "llvm.floor.f64";
else {
assert(0);
return;
}
- LLVMValueRef floor = lp_build_intrinsic(builder, intr, emit_data->dst_type,
+ LLVMValueRef floor = lp_build_intrinsic(ctx->ac.builder, intr, emit_data->dst_type,
&emit_data->args[0], 1,
LP_FUNC_ATTR_READNONE);
- emit_data->output[emit_data->chan] = LLVMBuildFSub(builder,
+ emit_data->output[emit_data->chan] = LLVMBuildFSub(ctx->ac.builder,
emit_data->args[0], floor, "");
}
static void emit_f2i(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildFPToSI(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildFPToSI(ctx->ac.builder,
emit_data->args[0], bld_base->int_bld.elem_type, "");
}
static void emit_f2u(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildFPToUI(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildFPToUI(ctx->ac.builder,
emit_data->args[0], bld_base->uint_bld.elem_type, "");
}
static void emit_i2f(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildSIToFP(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildSIToFP(ctx->ac.builder,
emit_data->args[0], bld_base->base.elem_type, "");
}
static void emit_u2f(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
- emit_data->output[emit_data->chan] = LLVMBuildUIToFP(builder,
+ struct si_shader_context *ctx = si_shader_context(bld_base);
+ emit_data->output[emit_data->chan] = LLVMBuildUIToFP(ctx->ac.builder,
emit_data->args[0], bld_base->base.elem_type, "");
}
static void
build_tgsi_intrinsic_nomem(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- struct lp_build_context *base = &bld_base->base;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
emit_data->output[emit_data->chan] =
- lp_build_intrinsic(base->gallivm->builder, action->intr_name,
+ lp_build_intrinsic(ctx->ac.builder, action->intr_name,
emit_data->dst_type, emit_data->args,
emit_data->arg_count, LP_FUNC_ATTR_READNONE);
}
static void emit_bfi(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct gallivm_state *gallivm = bld_base->base.gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef bfi_args[3];
LLVMValueRef bfi_sm5;
LLVMValueRef cond;
// Calculate the bitmask: (((1 << src3) - 1) << src2
bfi_args[0] = LLVMBuildShl(builder,
LLVMBuildSub(builder,
LLVMBuildShl(builder,
ctx->i32_1,
emit_data->args[3], ""),
@@ -523,67 +520,64 @@ static void emit_bfi(const struct lp_build_tgsi_action *action,
lp_build_const_int32(gallivm, 32), "");
emit_data->output[emit_data->chan] =
LLVMBuildSelect(builder, cond, emit_data->args[1], bfi_sm5, "");
}
static void emit_bfe(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef bfe_sm5;
LLVMValueRef cond;
bfe_sm5 = ac_build_bfe(&ctx->ac, emit_data->args[0],
emit_data->args[1], emit_data->args[2],
emit_data->info->opcode == TGSI_OPCODE_IBFE);
/* Correct for GLSL semantics. */
- cond = LLVMBuildICmp(builder, LLVMIntUGE, emit_data->args[2],
+ cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntUGE, emit_data->args[2],
LLVMConstInt(ctx->i32, 32, 0), "");
emit_data->output[emit_data->chan] =
- LLVMBuildSelect(builder, cond, emit_data->args[0], bfe_sm5, "");
+ LLVMBuildSelect(ctx->ac.builder, cond, emit_data->args[0], bfe_sm5, "");
}
/* this is ffs in C */
static void emit_lsb(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct gallivm_state *gallivm = bld_base->base.gallivm;
- LLVMBuilderRef builder = gallivm->builder;
LLVMValueRef args[2] = {
emit_data->args[0],
/* The value of 1 means that ffs(x=0) = undef, so LLVM won't
* add special code to check for x=0. The reason is that
* the LLVM behavior for x=0 is different from what we
* need here. However, LLVM also assumes that ffs(x) is
* in [0, 31], but GLSL expects that ffs(0) = -1, so
* a conditional assignment to handle 0 is still required.
*/
LLVMConstInt(LLVMInt1TypeInContext(gallivm->context), 1, 0)
};
LLVMValueRef lsb =
- lp_build_intrinsic(gallivm->builder, "llvm.cttz.i32",
+ lp_build_intrinsic(ctx->ac.builder, "llvm.cttz.i32",
emit_data->dst_type, args, ARRAY_SIZE(args),
LP_FUNC_ATTR_READNONE);
/* TODO: We need an intrinsic to skip this conditional. */
/* Check for zero: */
emit_data->output[emit_data->chan] =
- LLVMBuildSelect(builder,
- LLVMBuildICmp(builder, LLVMIntEQ, args[0],
+ LLVMBuildSelect(ctx->ac.builder,
+ LLVMBuildICmp(ctx->ac.builder, LLVMIntEQ, args[0],
ctx->i32_0, ""),
lp_build_const_int32(gallivm, -1), lsb, "");
}
/* Find the last bit set. */
static void emit_umsb(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
@@ -600,34 +594,34 @@ static void emit_imsb(const struct lp_build_tgsi_action *action,
struct si_shader_context *ctx = si_shader_context(bld_base);
emit_data->output[emit_data->chan] =
ac_build_imsb(&ctx->ac, emit_data->args[0],
emit_data->dst_type);
}
static void emit_iabs(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
emit_data->output[emit_data->chan] =
lp_build_emit_llvm_binary(bld_base, TGSI_OPCODE_IMAX,
emit_data->args[0],
- LLVMBuildNeg(builder,
+ LLVMBuildNeg(ctx->ac.builder,
emit_data->args[0], ""));
}
static void emit_minmax_int(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMIntPredicate op;
switch (emit_data->info->opcode) {
default:
assert(0);
case TGSI_OPCODE_IMAX:
case TGSI_OPCODE_I64MAX:
op = LLVMIntSGT;
break;
case TGSI_OPCODE_IMIN:
@@ -638,22 +632,22 @@ static void emit_minmax_int(const struct lp_build_tgsi_action *action,
case TGSI_OPCODE_U64MAX:
op = LLVMIntUGT;
break;
case TGSI_OPCODE_UMIN:
case TGSI_OPCODE_U64MIN:
op = LLVMIntULT;
break;
}
emit_data->output[emit_data->chan] =
- LLVMBuildSelect(builder,
- LLVMBuildICmp(builder, op, emit_data->args[0],
+ LLVMBuildSelect(ctx->ac.builder,
+ LLVMBuildICmp(ctx->ac.builder, op, emit_data->args[0],
emit_data->args[1], ""),
emit_data->args[0],
emit_data->args[1], "");
}
static void pk2h_fetch_args(struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
emit_data->args[0] = lp_build_emit_fetch(bld_base, emit_data->inst,
0, TGSI_CHAN_X);
@@ -680,47 +674,46 @@ static void up2h_fetch_args(struct lp_build_tgsi_context *bld_base,
{
emit_data->args[0] = lp_build_emit_fetch(bld_base, emit_data->inst,
0, TGSI_CHAN_X);
}
static void emit_up2h(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
LLVMContextRef context = bld_base->base.gallivm->context;
struct lp_build_context *uint_bld = &bld_base->uint_bld;
LLVMTypeRef i16;
LLVMValueRef const16, input, val;
unsigned i;
i16 = LLVMInt16TypeInContext(context);
const16 = lp_build_const_int32(uint_bld->gallivm, 16);
input = emit_data->args[0];
for (i = 0; i < 2; i++) {
- val = i == 1 ? LLVMBuildLShr(builder, input, const16, "") : input;
- val = LLVMBuildTrunc(builder, val, i16, "");
+ val = i == 1 ? LLVMBuildLShr(ctx->ac.builder, input, const16, "") : input;
+ val = LLVMBuildTrunc(ctx->ac.builder, val, i16, "");
val = ac_to_float(&ctx->ac, val);
- emit_data->output[i] = LLVMBuildFPExt(builder, val, ctx->f32, "");
+ emit_data->output[i] = LLVMBuildFPExt(ctx->ac.builder, val, ctx->f32, "");
}
}
static void emit_fdiv(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
emit_data->output[emit_data->chan] =
- LLVMBuildFDiv(ctx->gallivm.builder,
+ LLVMBuildFDiv(ctx->ac.builder,
emit_data->args[0], emit_data->args[1], "");
/* Use v_rcp_f32 instead of precise division. */
if (!LLVMIsConstant(emit_data->output[emit_data->chan]))
LLVMSetMetadata(emit_data->output[emit_data->chan],
ctx->fpmath_md_kind, ctx->fpmath_md_2p5_ulp);
}
/* 1/sqrt is translated to rsq for f32 if fp32 denormals are not enabled in
* the target machine. f64 needs global unsafe math flags to get rsq. */
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c b/src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c
index 887475b..32c44ba 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_mem.c
@@ -39,22 +39,21 @@ static const struct lp_build_tgsi_action tex_action;
/**
* Given a v8i32 resource descriptor for a buffer, extract the size of the
* buffer in number of elements and return it as an i32.
*/
static LLVMValueRef get_buffer_size(
struct lp_build_tgsi_context *bld_base,
LLVMValueRef descriptor)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef size =
LLVMBuildExtractElement(builder, descriptor,
LLVMConstInt(ctx->i32, 2, 0), "");
if (ctx->screen->b.chip_class == VI) {
/* On VI, the descriptor contains the size in bytes,
* but TXQ must return the size in elements.
* The stride is always non-zero for resources using TXQ.
*/
LLVMValueRef stride =
@@ -122,36 +121,35 @@ static bool tgsi_is_array_image(unsigned target)
* program termination) in this case, but it doesn't cost much to be a bit
* nicer: disabling DCC in the shader still leads to undefined results but
* avoids the lockup.
*/
static LLVMValueRef force_dcc_off(struct si_shader_context *ctx,
LLVMValueRef rsrc)
{
if (ctx->screen->b.chip_class <= CIK) {
return rsrc;
} else {
- LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef i32_6 = LLVMConstInt(ctx->i32, 6, 0);
LLVMValueRef i32_C = LLVMConstInt(ctx->i32, C_008F28_COMPRESSION_EN, 0);
LLVMValueRef tmp;
- tmp = LLVMBuildExtractElement(builder, rsrc, i32_6, "");
- tmp = LLVMBuildAnd(builder, tmp, i32_C, "");
- return LLVMBuildInsertElement(builder, rsrc, tmp, i32_6, "");
+ tmp = LLVMBuildExtractElement(ctx->ac.builder, rsrc, i32_6, "");
+ tmp = LLVMBuildAnd(ctx->ac.builder, tmp, i32_C, "");
+ return LLVMBuildInsertElement(ctx->ac.builder, rsrc, tmp, i32_6, "");
}
}
LLVMValueRef si_load_image_desc(struct si_shader_context *ctx,
LLVMValueRef list, LLVMValueRef index,
enum ac_descriptor_type desc_type, bool dcc_off)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef rsrc;
if (desc_type == AC_DESC_BUFFER) {
index = LLVMBuildMul(builder, index,
LLVMConstInt(ctx->i32, 2, 0), "");
index = LLVMBuildAdd(builder, index,
ctx->i32_1, "");
list = LLVMBuildPointerCast(builder, list,
si_const_array(ctx->v4i32, 0), "");
} else {
@@ -196,57 +194,54 @@ image_fetch_rsrc(
* If a shader performs an image load, store, or atomic
* operation using an image variable declared as an array,
* and if the index used to select an individual element is
* negative or greater than or equal to the size of the
* array, the results of the operation are undefined but may
* not lead to termination.
*/
index = si_get_bounded_indirect_index(ctx, &image->Indirect,
image->Register.Index,
ctx->num_images);
- index = LLVMBuildSub(ctx->gallivm.builder,
+ index = LLVMBuildSub(ctx->ac.builder,
LLVMConstInt(ctx->i32, SI_NUM_IMAGES - 1, 0),
index, "");
}
if (image->Register.File != TGSI_FILE_IMAGE) {
/* Bindless descriptors are accessible from a different pair of
* user SGPR indices.
*/
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
-
rsrc_ptr = LLVMGetParam(ctx->main_fn,
ctx->param_bindless_samplers_and_images);
index = lp_build_emit_fetch_src(bld_base, image,
TGSI_TYPE_UNSIGNED, 0);
/* For simplicity, bindless image descriptors use fixed
* 16-dword slots for now.
*/
- index = LLVMBuildMul(builder, index,
+ index = LLVMBuildMul(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, 2, 0), "");
}
*rsrc = si_load_image_desc(ctx, rsrc_ptr, index,
target == TGSI_TEXTURE_BUFFER ? AC_DESC_BUFFER : AC_DESC_IMAGE,
dcc_off);
}
static LLVMValueRef image_fetch_coords(
struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_instruction *inst,
unsigned src, LLVMValueRef desc)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
unsigned target = inst->Memory.Texture;
unsigned num_coords = tgsi_util_get_texture_coord_dim(target);
LLVMValueRef coords[4];
LLVMValueRef tmp;
int chan;
for (chan = 0; chan < num_coords; ++chan) {
tmp = lp_build_emit_fetch(bld_base, inst, src, chan);
tmp = ac_to_integer(&ctx->ac, tmp);
coords[chan] = tmp;
@@ -452,58 +447,56 @@ static void load_emit_buffer(struct si_shader_context *ctx,
args[2], NULL, 0,
LLVMConstIntGetZExtValue(args[3]),
LLVMConstIntGetZExtValue(args[4]),
can_speculate, allow_smem);
}
static LLVMValueRef get_memory_ptr(struct si_shader_context *ctx,
const struct tgsi_full_instruction *inst,
LLVMTypeRef type, int arg)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef offset, ptr;
int addr_space;
offset = lp_build_emit_fetch(&ctx->bld_base, inst, arg, 0);
offset = ac_to_integer(&ctx->ac, offset);
ptr = ctx->shared_memory;
ptr = LLVMBuildGEP(builder, ptr, &offset, 1, "");
addr_space = LLVMGetPointerAddressSpace(LLVMTypeOf(ptr));
ptr = LLVMBuildBitCast(builder, ptr, LLVMPointerType(type, addr_space), "");
return ptr;
}
static void load_emit_memory(
struct si_shader_context *ctx,
struct lp_build_emit_data *emit_data)
{
const struct tgsi_full_instruction *inst = emit_data->inst;
struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
unsigned writemask = inst->Dst[0].Register.WriteMask;
LLVMValueRef channels[4], ptr, derived_ptr, index;
int chan;
ptr = get_memory_ptr(ctx, inst, ctx->f32, 1);
for (chan = 0; chan < 4; ++chan) {
if (!(writemask & (1 << chan))) {
channels[chan] = LLVMGetUndef(ctx->f32);
continue;
}
index = LLVMConstInt(ctx->i32, chan, 0);
- derived_ptr = LLVMBuildGEP(builder, ptr, &index, 1, "");
- channels[chan] = LLVMBuildLoad(builder, derived_ptr, "");
+ derived_ptr = LLVMBuildGEP(ctx->ac.builder, ptr, &index, 1, "");
+ channels[chan] = LLVMBuildLoad(ctx->ac.builder, derived_ptr, "");
}
emit_data->output[emit_data->chan] = lp_build_gather_values(gallivm, channels, 4);
}
/**
* Return true if the memory accessed by a LOAD or STORE instruction is
* read-only or write-only, respectively.
*
* \param shader_buffers_reverse_access_mask
* For LOAD, set this to (store | atomic) slot usage in the shader.
@@ -563,22 +556,21 @@ static bool is_oneway_access_only(const struct tgsi_full_instruction *inst,
}
return false;
}
static void load_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
const struct tgsi_shader_info *info = &ctx->shader->selector->info;
char intrinsic_name[64];
bool can_speculate = false;
if (inst->Src[0].Register.File == TGSI_FILE_MEMORY) {
load_emit_memory(ctx, emit_data);
return;
}
@@ -687,22 +679,21 @@ static void store_fetch_args(
}
}
}
static void store_emit_buffer(
struct si_shader_context *ctx,
struct lp_build_emit_data *emit_data,
bool writeonly_memory)
{
const struct tgsi_full_instruction *inst = emit_data->inst;
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef base_data = emit_data->args[0];
LLVMValueRef base_offset = emit_data->args[3];
unsigned writemask = inst->Dst[0].Register.WriteMask;
while (writemask) {
int start, count;
const char *intrinsic_name;
LLVMValueRef data;
LLVMValueRef offset;
LLVMValueRef tmp;
@@ -759,22 +750,21 @@ static void store_emit_buffer(
emit_data->args, emit_data->arg_count,
get_store_intr_attribs(writeonly_memory));
}
}
static void store_emit_memory(
struct si_shader_context *ctx,
struct lp_build_emit_data *emit_data)
{
const struct tgsi_full_instruction *inst = emit_data->inst;
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
unsigned writemask = inst->Dst[0].Register.WriteMask;
LLVMValueRef ptr, derived_ptr, data, index;
int chan;
ptr = get_memory_ptr(ctx, inst, ctx->f32, 0);
for (chan = 0; chan < 4; ++chan) {
if (!(writemask & (1 << chan))) {
continue;
}
@@ -784,22 +774,21 @@ static void store_emit_memory(
LLVMBuildStore(builder, data, derived_ptr);
}
}
static void store_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
const struct tgsi_shader_info *info = &ctx->shader->selector->info;
unsigned target = inst->Memory.Texture;
char intrinsic_name[64];
bool writeonly_memory = false;
if (inst->Dst[0].Register.File == TGSI_FILE_MEMORY) {
store_emit_memory(ctx, emit_data);
return;
}
@@ -891,22 +880,21 @@ static void atomic_fetch_args(
emit_data->args[emit_data->arg_count++] = coords;
emit_data->args[emit_data->arg_count++] = rsrc;
image_append_args(ctx, emit_data, target, true, false);
}
}
}
static void atomic_emit_memory(struct si_shader_context *ctx,
struct lp_build_emit_data *emit_data) {
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
LLVMValueRef ptr, result, arg;
ptr = get_memory_ptr(ctx, inst, ctx->i32, 1);
arg = lp_build_emit_fetch(&ctx->bld_base, inst, 2, 0);
arg = ac_to_integer(&ctx->ac, arg);
if (inst->Instruction.Opcode == TGSI_OPCODE_ATOMCAS) {
LLVMValueRef new_data;
@@ -962,22 +950,21 @@ static void atomic_emit_memory(struct si_shader_context *ctx,
}
emit_data->output[emit_data->chan] = LLVMBuildBitCast(builder, result, emit_data->dst_type, "");
}
static void atomic_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
const struct tgsi_full_instruction * inst = emit_data->inst;
char intrinsic_name[40];
LLVMValueRef tmp;
if (inst->Src[0].Register.File == TGSI_FILE_MEMORY) {
atomic_emit_memory(ctx, emit_data);
return;
}
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER ||
@@ -1032,21 +1019,21 @@ static void set_tex_fetch_args(struct si_shader_context *ctx,
args.da = tgsi_is_array_sampler(target);
/* Ugly, but we seem to have no other choice right now. */
STATIC_ASSERT(sizeof(args) <= sizeof(emit_data->args));
memcpy(emit_data->args, &args, sizeof(args));
}
static LLVMValueRef fix_resinfo(struct si_shader_context *ctx,
unsigned target, LLVMValueRef out)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
/* 1D textures are allocated and used as 2D on GFX9. */
if (ctx->screen->b.chip_class >= GFX9 &&
(target == TGSI_TEXTURE_1D_ARRAY ||
target == TGSI_TEXTURE_SHADOW1D_ARRAY)) {
LLVMValueRef layers =
LLVMBuildExtractElement(builder, out,
LLVMConstInt(ctx->i32, 2, 0), "");
out = LLVMBuildInsertElement(builder, out, layers,
ctx->i32_1, "");
@@ -1098,22 +1085,21 @@ static void resq_fetch_args(
0xf);
}
}
static void resq_emit(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
const struct tgsi_full_instruction *inst = emit_data->inst;
LLVMValueRef out;
if (inst->Src[0].Register.File == TGSI_FILE_BUFFER) {
out = LLVMBuildExtractElement(builder, emit_data->args[0],
LLVMConstInt(ctx->i32, 2, 0), "");
} else if (inst->Memory.Texture == TGSI_TEXTURE_BUFFER) {
out = get_buffer_size(bld_base, emit_data->args[0]);
} else {
struct ac_image_args args;
@@ -1128,22 +1114,21 @@ static void resq_emit(
emit_data->output[emit_data->chan] = out;
}
/**
* Load an image view, fmask view. or sampler state descriptor.
*/
LLVMValueRef si_load_sampler_desc(struct si_shader_context *ctx,
LLVMValueRef list, LLVMValueRef index,
enum ac_descriptor_type type)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
switch (type) {
case AC_DESC_IMAGE:
/* The image is at [0:7]. */
index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, 2, 0), "");
break;
case AC_DESC_BUFFER:
/* The buffer is in [4:7]. */
index = LLVMBuildMul(builder, index, LLVMConstInt(ctx->i32, 4, 0), "");
index = LLVMBuildAdd(builder, index, ctx->i32_1, "");
@@ -1174,32 +1159,31 @@ LLVMValueRef si_load_sampler_desc(struct si_shader_context *ctx,
* filtering manually. The driver sets img7 to a mask clearing
* MAX_ANISO_RATIO if BASE_LEVEL == LAST_LEVEL. The shader must do:
* s_and_b32 samp0, samp0, img7
*
* VI:
* The ANISO_OVERRIDE sampler field enables this fix in TA.
*/
static LLVMValueRef sici_fix_sampler_aniso(struct si_shader_context *ctx,
LLVMValueRef res, LLVMValueRef samp)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
LLVMValueRef img7, samp0;
if (ctx->screen->b.chip_class >= VI)
return samp;
- img7 = LLVMBuildExtractElement(builder, res,
+ img7 = LLVMBuildExtractElement(ctx->ac.builder, res,
LLVMConstInt(ctx->i32, 7, 0), "");
- samp0 = LLVMBuildExtractElement(builder, samp,
+ samp0 = LLVMBuildExtractElement(ctx->ac.builder, samp,
ctx->i32_0, "");
- samp0 = LLVMBuildAnd(builder, samp0, img7, "");
- return LLVMBuildInsertElement(builder, samp, samp0,
+ samp0 = LLVMBuildAnd(ctx->ac.builder, samp0, img7, "");
+ return LLVMBuildInsertElement(ctx->ac.builder, samp, samp0,
ctx->i32_0, "");
}
static void tex_fetch_ptrs(
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data,
LLVMValueRef *res_ptr, LLVMValueRef *samp_ptr, LLVMValueRef *fmask_ptr)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef list = LLVMGetParam(ctx->main_fn, ctx->param_samplers_and_images);
@@ -1210,21 +1194,21 @@ static void tex_fetch_ptrs(
LLVMValueRef index;
sampler_src = emit_data->inst->Instruction.NumSrcRegs - 1;
reg = &emit_data->inst->Src[sampler_src];
if (reg->Register.Indirect) {
index = si_get_bounded_indirect_index(ctx,
®->Indirect,
reg->Register.Index,
ctx->num_samplers);
- index = LLVMBuildAdd(ctx->gallivm.builder, index,
+ index = LLVMBuildAdd(ctx->ac.builder, index,
LLVMConstInt(ctx->i32, SI_NUM_IMAGES / 2, 0), "");
} else {
index = LLVMConstInt(ctx->i32,
si_get_sampler_slot(reg->Register.Index), 0);
}
if (reg->Register.File != TGSI_FILE_SAMPLER) {
/* Bindless descriptors are accessible from a different pair of
* user SGPR indices.
*/
@@ -1303,21 +1287,20 @@ static void txq_emit(const struct lp_build_tgsi_action *action,
LLVMValueRef result = ac_build_image_opcode(&ctx->ac, &args);
emit_data->output[emit_data->chan] = fix_resinfo(ctx, target, result);
}
static void tex_fetch_args(
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_instruction *inst = emit_data->inst;
unsigned opcode = inst->Instruction.Opcode;
unsigned target = inst->Texture.Texture;
LLVMValueRef coords[5], derivs[6];
LLVMValueRef address[16];
unsigned num_coords = tgsi_util_get_texture_coord_dim(target);
int ref_pos = tgsi_util_get_shadow_ref_src_index(target);
unsigned count = 0;
unsigned chan;
unsigned num_deriv_channels = 0;
@@ -1359,29 +1342,29 @@ static void tex_fetch_args(
/* The offsets are six-bit signed integers packed like this:
* X=[5:0], Y=[13:8], and Z=[21:16].
*/
LLVMValueRef offset[3], pack;
assert(inst->Texture.NumOffsets == 1);
for (chan = 0; chan < 3; chan++) {
offset[chan] = lp_build_emit_fetch_texoffset(bld_base,
emit_data->inst, 0, chan);
- offset[chan] = LLVMBuildAnd(gallivm->builder, offset[chan],
+ offset[chan] = LLVMBuildAnd(ctx->ac.builder, offset[chan],
LLVMConstInt(ctx->i32, 0x3f, 0), "");
if (chan)
- offset[chan] = LLVMBuildShl(gallivm->builder, offset[chan],
+ offset[chan] = LLVMBuildShl(ctx->ac.builder, offset[chan],
LLVMConstInt(ctx->i32, chan*8, 0), "");
}
- pack = LLVMBuildOr(gallivm->builder, offset[0], offset[1], "");
- pack = LLVMBuildOr(gallivm->builder, pack, offset[2], "");
+ pack = LLVMBuildOr(ctx->ac.builder, offset[0], offset[1], "");
+ pack = LLVMBuildOr(ctx->ac.builder, pack, offset[2], "");
address[count++] = pack;
}
/* Pack LOD bias value */
if (opcode == TGSI_OPCODE_TXB)
address[count++] = coords[3];
if (opcode == TGSI_OPCODE_TXB2)
address[count++] = lp_build_emit_fetch(bld_base, inst, 1, TGSI_CHAN_X);
/* Pack depth comparison value */
@@ -1402,27 +1385,27 @@ static void tex_fetch_args(
* depth texture, then D_t and D_ref are clamped to the
* range [0, 1]; otherwise no clamping is performed."
*
* TC-compatible HTILE promotes Z16 and Z24 to Z32_FLOAT,
* so the depth comparison value isn't clamped for Z16 and
* Z24 anymore. Do it manually here.
*/
if (ctx->screen->b.chip_class >= VI) {
LLVMValueRef upgraded;
LLVMValueRef clamped;
- upgraded = LLVMBuildExtractElement(gallivm->builder, samp_ptr,
+ upgraded = LLVMBuildExtractElement(ctx->ac.builder, samp_ptr,
LLVMConstInt(ctx->i32, 3, false), "");
- upgraded = LLVMBuildLShr(gallivm->builder, upgraded,
+ upgraded = LLVMBuildLShr(ctx->ac.builder, upgraded,
LLVMConstInt(ctx->i32, 29, false), "");
- upgraded = LLVMBuildTrunc(gallivm->builder, upgraded, ctx->i1, "");
+ upgraded = LLVMBuildTrunc(ctx->ac.builder, upgraded, ctx->i1, "");
clamped = ac_build_clamp(&ctx->ac, z);
- z = LLVMBuildSelect(gallivm->builder, upgraded, clamped, z, "");
+ z = LLVMBuildSelect(ctx->ac.builder, upgraded, clamped, z, "");
}
address[count++] = z;
}
/* Pack user derivatives */
if (opcode == TGSI_OPCODE_TXD) {
int param, num_src_deriv_channels, num_dst_deriv_channels;
switch (target) {
@@ -1582,53 +1565,53 @@ static void tex_fetch_args(
target, fmask_ptr, NULL,
txf_address, txf_count, 0xf);
build_tex_intrinsic(&tex_action, bld_base, &txf_emit_data);
/* Initialize some constants. */
LLVMValueRef four = LLVMConstInt(ctx->i32, 4, 0);
LLVMValueRef F = LLVMConstInt(ctx->i32, 0xF, 0);
/* Apply the formula. */
LLVMValueRef fmask =
- LLVMBuildExtractElement(gallivm->builder,
+ LLVMBuildExtractElement(ctx->ac.builder,
txf_emit_data.output[0],
ctx->i32_0, "");
unsigned sample_chan = txf_count; /* the sample index is last */
LLVMValueRef sample_index4 =
- LLVMBuildMul(gallivm->builder, address[sample_chan], four, "");
+ LLVMBuildMul(ctx->ac.builder, address[sample_chan], four, "");
LLVMValueRef shifted_fmask =
- LLVMBuildLShr(gallivm->builder, fmask, sample_index4, "");
+ LLVMBuildLShr(ctx->ac.builder, fmask, sample_index4, "");
LLVMValueRef final_sample =
- LLVMBuildAnd(gallivm->builder, shifted_fmask, F, "");
+ LLVMBuildAnd(ctx->ac.builder, shifted_fmask, F, "");
/* Don't rewrite the sample index if WORD1.DATA_FORMAT of the FMASK
* resource descriptor is 0 (invalid),
*/
LLVMValueRef fmask_desc =
- LLVMBuildBitCast(gallivm->builder, fmask_ptr,
+ LLVMBuildBitCast(ctx->ac.builder, fmask_ptr,
ctx->v8i32, "");
LLVMValueRef fmask_word1 =
- LLVMBuildExtractElement(gallivm->builder, fmask_desc,
+ LLVMBuildExtractElement(ctx->ac.builder, fmask_desc,
ctx->i32_1, "");
LLVMValueRef word1_is_nonzero =
- LLVMBuildICmp(gallivm->builder, LLVMIntNE,
+ LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
fmask_word1, ctx->i32_0, "");
/* Replace the MSAA sample index. */
address[sample_chan] =
- LLVMBuildSelect(gallivm->builder, word1_is_nonzero,
+ LLVMBuildSelect(ctx->ac.builder, word1_is_nonzero,
final_sample, address[sample_chan], "");
}
if (opcode == TGSI_OPCODE_TXF ||
opcode == TGSI_OPCODE_TXF_LZ) {
/* add tex offsets */
if (inst->Texture.NumOffsets) {
struct lp_build_context *uint_bld = &bld_base->uint_bld;
const struct tgsi_texture_offset *off = inst->TexOffsets;
@@ -1706,21 +1689,21 @@ static void tex_fetch_args(
* precision in 32-bit data formats, so it needs to be applied dynamically at
* runtime. In this case, return an i1 value that indicates whether the
* descriptor was overridden (and hence a fixup of the sampler result is needed).
*/
static LLVMValueRef
si_lower_gather4_integer(struct si_shader_context *ctx,
struct ac_image_args *args,
unsigned target,
enum tgsi_return_type return_type)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef coord = args->addr;
LLVMValueRef half_texel[2];
/* Texture coordinates start after:
* {offset, bias, z-compare, derivatives}
* Only the offset and z-compare can occur here.
*/
unsigned coord_vgpr_index = (int)args->offset + (int)args->compare;
int c;
assert(return_type == TGSI_RETURN_TYPE_SINT ||
@@ -1808,21 +1791,21 @@ si_lower_gather4_integer(struct si_shader_context *ctx,
/* The second half of the cube texture 8_8_8_8 integer workaround: adjust the
* result after the gather operation.
*/
static LLVMValueRef
si_fix_gather4_integer_result(struct si_shader_context *ctx,
LLVMValueRef result,
enum tgsi_return_type return_type,
LLVMValueRef wa)
{
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
assert(return_type == TGSI_RETURN_TYPE_SINT ||
return_type == TGSI_RETURN_TYPE_UINT);
for (unsigned chan = 0; chan < 4; ++chan) {
LLVMValueRef chanv = LLVMConstInt(ctx->i32, chan, false);
LLVMValueRef value;
LLVMValueRef wa_value;
value = LLVMBuildExtractElement(builder, result, chanv, "");
@@ -1937,38 +1920,37 @@ static void build_tex_intrinsic(const struct lp_build_tgsi_action *action,
emit_data->output[emit_data->chan] = result;
}
static void si_llvm_emit_txqs(
const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef res, samples;
LLVMValueRef res_ptr, samp_ptr, fmask_ptr = NULL;
tex_fetch_ptrs(bld_base, emit_data, &res_ptr, &samp_ptr, &fmask_ptr);
/* Read the samples from the descriptor directly. */
- res = LLVMBuildBitCast(builder, res_ptr, ctx->v8i32, "");
+ res = LLVMBuildBitCast(ctx->ac.builder, res_ptr, ctx->v8i32, "");
samples = LLVMBuildExtractElement(
builder, res,
LLVMConstInt(ctx->i32, 3, 0), "");
- samples = LLVMBuildLShr(builder, samples,
+ samples = LLVMBuildLShr(ctx->ac.builder, samples,
LLVMConstInt(ctx->i32, 16, 0), "");
- samples = LLVMBuildAnd(builder, samples,
+ samples = LLVMBuildAnd(ctx->ac.builder, samples,
LLVMConstInt(ctx->i32, 0xf, 0), "");
- samples = LLVMBuildShl(builder, ctx->i32_1,
+ samples = LLVMBuildShl(ctx->ac.builder, ctx->i32_1,
samples, "");
emit_data->output[emit_data->chan] = samples;
}
static const struct lp_build_tgsi_action tex_action = {
.fetch_args = tex_fetch_args,
.emit = build_tex_intrinsic,
};
diff --git a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
index aee8cdd..2ecb112 100644
--- a/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
+++ b/src/gallium/drivers/radeonsi/si_shader_tgsi_setup.c
@@ -180,39 +180,38 @@ LLVMTypeRef tgsi2llvmtype(struct lp_build_tgsi_context *bld_base,
case TGSI_TYPE_FLOAT:
return LLVMFloatTypeInContext(ctx);
default: break;
}
return 0;
}
LLVMValueRef bitcast(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type, LLVMValueRef value)
{
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMTypeRef dst_type = tgsi2llvmtype(bld_base, type);
if (dst_type)
- return LLVMBuildBitCast(builder, value, dst_type, "");
+ return LLVMBuildBitCast(ctx->ac.builder, value, dst_type, "");
else
return value;
}
/**
* Return a value that is equal to the given i32 \p index if it lies in [0,num)
* or an undefined value in the same interval otherwise.
*/
LLVMValueRef si_llvm_bound_index(struct si_shader_context *ctx,
LLVMValueRef index,
unsigned num)
{
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef c_max = LLVMConstInt(ctx->i32, num - 1, 0);
LLVMValueRef cc;
if (util_is_power_of_two(num)) {
index = LLVMBuildAnd(builder, index, c_max, "");
} else {
/* In theory, this MAX pattern should result in code that is
* as good as the bit-wise AND above.
*
* In practice, LLVM generates worse code (at the time of
@@ -264,30 +263,31 @@ push_flow(struct si_shader_context *ctx)
return flow;
}
static LLVMValueRef emit_swizzle(struct lp_build_tgsi_context *bld_base,
LLVMValueRef value,
unsigned swizzle_x,
unsigned swizzle_y,
unsigned swizzle_z,
unsigned swizzle_w)
{
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef swizzles[4];
LLVMTypeRef i32t =
LLVMInt32TypeInContext(bld_base->base.gallivm->context);
swizzles[0] = LLVMConstInt(i32t, swizzle_x, 0);
swizzles[1] = LLVMConstInt(i32t, swizzle_y, 0);
swizzles[2] = LLVMConstInt(i32t, swizzle_z, 0);
swizzles[3] = LLVMConstInt(i32t, swizzle_w, 0);
- return LLVMBuildShuffleVector(bld_base->base.gallivm->builder,
+ return LLVMBuildShuffleVector(ctx->ac.builder,
value,
LLVMGetUndef(LLVMTypeOf(value)),
LLVMConstVector(swizzles, 4), "");
}
/**
* Return the description of the array covering the given temporary register
* index.
*/
static unsigned
@@ -340,22 +340,21 @@ get_array_range(struct lp_build_tgsi_context *bld_base,
*/
static LLVMValueRef
get_pointer_into_array(struct si_shader_context *ctx,
unsigned file,
unsigned swizzle,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
unsigned array_id;
struct tgsi_array_info *array;
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef idxs[2];
LLVMValueRef index;
LLVMValueRef alloca;
if (file != TGSI_FILE_TEMPORARY)
return NULL;
array_id = get_temp_array_id(&ctx->bld_base, reg_index, reg_indirect);
if (!array_id)
return NULL;
@@ -390,83 +389,78 @@ get_pointer_into_array(struct si_shader_context *ctx,
builder, index,
LLVMConstInt(ctx->i32, util_bitcount(array->writemask), 0),
"");
index = LLVMBuildAdd(
builder, index,
LLVMConstInt(ctx->i32,
util_bitcount(array->writemask & ((1 << swizzle) - 1)), 0),
"");
idxs[0] = ctx->i32_0;
idxs[1] = index;
- return LLVMBuildGEP(builder, alloca, idxs, 2, "");
+ return LLVMBuildGEP(ctx->ac.builder, alloca, idxs, 2, "");
}
LLVMValueRef
si_llvm_emit_fetch_64bit(struct lp_build_tgsi_context *bld_base,
enum tgsi_opcode_type type,
LLVMValueRef ptr,
LLVMValueRef ptr2)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = bld_base->base.gallivm->builder;
LLVMValueRef result;
result = LLVMGetUndef(LLVMVectorType(LLVMIntTypeInContext(bld_base->base.gallivm->context, 32), bld_base->base.type.length * 2));
- result = LLVMBuildInsertElement(builder,
+ result = LLVMBuildInsertElement(ctx->ac.builder,
result,
ac_to_integer(&ctx->ac, ptr),
ctx->i32_0, "");
- result = LLVMBuildInsertElement(builder,
+ result = LLVMBuildInsertElement(ctx->ac.builder,
result,
ac_to_integer(&ctx->ac, ptr2),
ctx->i32_1, "");
return bitcast(bld_base, type, result);
}
static LLVMValueRef
emit_array_fetch(struct lp_build_tgsi_context *bld_base,
unsigned File, enum tgsi_opcode_type type,
struct tgsi_declaration_range range,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
-
- LLVMBuilderRef builder = ctx->gallivm.builder;
-
unsigned i, size = range.Last - range.First + 1;
LLVMTypeRef vec = LLVMVectorType(tgsi2llvmtype(bld_base, type), size);
LLVMValueRef result = LLVMGetUndef(vec);
struct tgsi_full_src_register tmp_reg = {};
tmp_reg.Register.File = File;
for (i = 0; i < size; ++i) {
tmp_reg.Register.Index = i + range.First;
LLVMValueRef temp = si_llvm_emit_fetch(bld_base, &tmp_reg, type, swizzle);
- result = LLVMBuildInsertElement(builder, result, temp,
+ result = LLVMBuildInsertElement(ctx->ac.builder, result, temp,
LLVMConstInt(ctx->i32, i, 0), "array_vector");
}
return result;
}
static LLVMValueRef
load_value_from_array(struct lp_build_tgsi_context *bld_base,
unsigned file,
enum tgsi_opcode_type type,
unsigned swizzle,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef ptr;
ptr = get_pointer_into_array(ctx, file, swizzle, reg_index, reg_indirect);
if (ptr) {
LLVMValueRef val = LLVMBuildLoad(builder, ptr, "");
if (tgsi_type_is_64bit(type)) {
LLVMValueRef ptr_hi, val_hi;
ptr_hi = LLVMBuildGEP(builder, ptr, &ctx->i32_1, 1, "");
val_hi = LLVMBuildLoad(builder, ptr_hi, "");
val = si_llvm_emit_fetch_64bit(bld_base, type, val, val_hi);
@@ -486,22 +480,21 @@ load_value_from_array(struct lp_build_tgsi_context *bld_base,
static void
store_value_to_array(struct lp_build_tgsi_context *bld_base,
LLVMValueRef value,
unsigned file,
unsigned chan_index,
unsigned reg_index,
const struct tgsi_ind_register *reg_indirect)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
- LLVMBuilderRef builder = gallivm->builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef ptr;
ptr = get_pointer_into_array(ctx, file, chan_index, reg_index, reg_indirect);
if (ptr) {
LLVMBuildStore(builder, value, ptr);
} else {
unsigned i, size;
struct tgsi_declaration_range range = get_array_range(bld_base, file, reg_index, reg_indirect);
LLVMValueRef index = si_get_indirect_index(ctx, reg_indirect, 1, reg_index - range.First);
LLVMValueRef array =
@@ -555,21 +548,21 @@ get_output_ptr(struct lp_build_tgsi_context *bld_base, unsigned index,
assert(index <= ctx->bld_base.info->file_max[TGSI_FILE_OUTPUT]);
return ctx->outputs[index][chan];
}
LLVMValueRef si_llvm_emit_fetch(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef result = NULL, ptr, ptr2;
if (swizzle == ~0) {
LLVMValueRef values[TGSI_NUM_CHANNELS];
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
values[chan] = si_llvm_emit_fetch(bld_base, reg, type, chan);
}
return lp_build_gather_values(&ctx->gallivm, values,
TGSI_NUM_CHANNELS);
@@ -653,21 +646,21 @@ LLVMValueRef si_llvm_emit_fetch(struct lp_build_tgsi_context *bld_base,
return bitcast(bld_base, type, result);
}
static LLVMValueRef fetch_system_value(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_src_register *reg,
enum tgsi_opcode_type type,
unsigned swizzle)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef cval = ctx->system_values[reg->Register.Index];
if (tgsi_type_is_64bit(type)) {
LLVMValueRef lo, hi;
assert(swizzle == 0 || swizzle == 2);
lo = LLVMBuildExtractElement(
builder, cval, LLVMConstInt(ctx->i32, swizzle, 0), "");
hi = LLVMBuildExtractElement(
@@ -683,21 +676,21 @@ static LLVMValueRef fetch_system_value(struct lp_build_tgsi_context *bld_base,
assert(swizzle == 0);
}
return bitcast(bld_base, type, cval);
}
static void emit_declaration(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_declaration *decl)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
unsigned first, last, i;
switch(decl->Declaration.File) {
case TGSI_FILE_ADDRESS:
{
unsigned idx;
for (idx = decl->Range.First; idx <= decl->Range.Last; idx++) {
unsigned chan;
for (chan = 0; chan < TGSI_NUM_CHANNELS; chan++) {
ctx->addrs[idx][chan] = lp_build_alloca_undef(
&ctx->gallivm,
@@ -865,39 +858,38 @@ static void emit_declaration(struct lp_build_tgsi_context *bld_base,
}
}
void si_llvm_emit_store(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_instruction *inst,
const struct tgsi_opcode_info *info,
unsigned index,
LLVMValueRef dst[4])
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
const struct tgsi_full_dst_register *reg = &inst->Dst[index];
- LLVMBuilderRef builder = ctx->gallivm.builder;
+ LLVMBuilderRef builder = ctx->ac.builder;
LLVMValueRef temp_ptr, temp_ptr2 = NULL;
bool is_vec_store = false;
enum tgsi_opcode_type dtype = tgsi_opcode_infer_dst_type(inst->Instruction.Opcode, index);
if (dst[0]) {
LLVMTypeKind k = LLVMGetTypeKind(LLVMTypeOf(dst[0]));
is_vec_store = (k == LLVMVectorTypeKind);
}
if (is_vec_store) {
LLVMValueRef values[4] = {};
uint32_t writemask = reg->Register.WriteMask;
while (writemask) {
unsigned chan = u_bit_scan(&writemask);
LLVMValueRef index = LLVMConstInt(ctx->i32, chan, 0);
- values[chan] = LLVMBuildExtractElement(gallivm->builder,
+ values[chan] = LLVMBuildExtractElement(ctx->ac.builder,
dst[0], index, "");
}
bld_base->emit_store(bld_base, inst, info, index, values);
return;
}
uint32_t writemask = reg->Register.WriteMask;
while (writemask) {
unsigned chan_index = u_bit_scan(&writemask);
LLVMValueRef value = dst[chan_index];
@@ -999,145 +991,137 @@ static void emit_default_branch(LLVMBuilderRef builder, LLVMBasicBlockRef target
{
if (!LLVMGetBasicBlockTerminator(LLVMGetInsertBlock(builder)))
LLVMBuildBr(builder, target);
}
static void bgnloop_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = push_flow(ctx);
flow->loop_entry_block = append_basic_block(ctx, "LOOP");
flow->next_block = append_basic_block(ctx, "ENDLOOP");
set_basicblock_name(flow->loop_entry_block, "loop", bld_base->pc);
- LLVMBuildBr(gallivm->builder, flow->loop_entry_block);
- LLVMPositionBuilderAtEnd(gallivm->builder, flow->loop_entry_block);
+ LLVMBuildBr(ctx->ac.builder, flow->loop_entry_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, flow->loop_entry_block);
}
static void brk_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = get_innermost_loop(ctx);
- LLVMBuildBr(gallivm->builder, flow->next_block);
+ LLVMBuildBr(ctx->ac.builder, flow->next_block);
}
static void cont_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = get_innermost_loop(ctx);
- LLVMBuildBr(gallivm->builder, flow->loop_entry_block);
+ LLVMBuildBr(ctx->ac.builder, flow->loop_entry_block);
}
static void else_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_branch = get_current_flow(ctx);
LLVMBasicBlockRef endif_block;
assert(!current_branch->loop_entry_block);
endif_block = append_basic_block(ctx, "ENDIF");
- emit_default_branch(gallivm->builder, endif_block);
+ emit_default_branch(ctx->ac.builder, endif_block);
- LLVMPositionBuilderAtEnd(gallivm->builder, current_branch->next_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, current_branch->next_block);
set_basicblock_name(current_branch->next_block, "else", bld_base->pc);
current_branch->next_block = endif_block;
}
static void endif_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_branch = get_current_flow(ctx);
assert(!current_branch->loop_entry_block);
- emit_default_branch(gallivm->builder, current_branch->next_block);
- LLVMPositionBuilderAtEnd(gallivm->builder, current_branch->next_block);
+ emit_default_branch(ctx->ac.builder, current_branch->next_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, current_branch->next_block);
set_basicblock_name(current_branch->next_block, "endif", bld_base->pc);
ctx->flow_depth--;
}
static void endloop_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *current_loop = get_current_flow(ctx);
assert(current_loop->loop_entry_block);
- emit_default_branch(gallivm->builder, current_loop->loop_entry_block);
+ emit_default_branch(ctx->ac.builder, current_loop->loop_entry_block);
- LLVMPositionBuilderAtEnd(gallivm->builder, current_loop->next_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, current_loop->next_block);
set_basicblock_name(current_loop->next_block, "endloop", bld_base->pc);
ctx->flow_depth--;
}
static void if_cond_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data,
LLVMValueRef cond)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = &ctx->gallivm;
struct si_llvm_flow *flow = push_flow(ctx);
LLVMBasicBlockRef if_block;
if_block = append_basic_block(ctx, "IF");
flow->next_block = append_basic_block(ctx, "ELSE");
set_basicblock_name(if_block, "if", bld_base->pc);
- LLVMBuildCondBr(gallivm->builder, cond, if_block, flow->next_block);
- LLVMPositionBuilderAtEnd(gallivm->builder, if_block);
+ LLVMBuildCondBr(ctx->ac.builder, cond, if_block, flow->next_block);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, if_block);
}
static void if_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
- struct gallivm_state *gallivm = bld_base->base.gallivm;
+ struct si_shader_context *ctx = si_shader_context(bld_base);
LLVMValueRef cond;
- cond = LLVMBuildFCmp(gallivm->builder, LLVMRealUNE,
+ cond = LLVMBuildFCmp(ctx->ac.builder, LLVMRealUNE,
emit_data->args[0],
bld_base->base.zero, "");
if_cond_emit(action, bld_base, emit_data, cond);
}
static void uif_emit(const struct lp_build_tgsi_action *action,
struct lp_build_tgsi_context *bld_base,
struct lp_build_emit_data *emit_data)
{
struct si_shader_context *ctx = si_shader_context(bld_base);
- struct gallivm_state *gallivm = bld_base->base.gallivm;
LLVMValueRef cond;
- cond = LLVMBuildICmp(gallivm->builder, LLVMIntNE,
+ cond = LLVMBuildICmp(ctx->ac.builder, LLVMIntNE,
ac_to_integer(&ctx->ac, emit_data->args[0]), ctx->i32_0, "");
if_cond_emit(action, bld_base, emit_data, cond);
}
static void emit_immediate(struct lp_build_tgsi_context *bld_base,
const struct tgsi_full_immediate *imm)
{
unsigned i;
struct si_shader_context *ctx = si_shader_context(bld_base);
@@ -1327,21 +1311,21 @@ void si_llvm_create_func(struct si_shader_context *ctx,
num_return_elems, true);
else
ret_type = LLVMVoidTypeInContext(ctx->gallivm.context);
/* Setup the function */
ctx->return_type = ret_type;
main_fn_type = LLVMFunctionType(ret_type, ParamTypes, ParamCount, 0);
ctx->main_fn = LLVMAddFunction(ctx->gallivm.module, name, main_fn_type);
main_fn_body = LLVMAppendBasicBlockInContext(ctx->gallivm.context,
ctx->main_fn, "main_body");
- LLVMPositionBuilderAtEnd(ctx->gallivm.builder, main_fn_body);
+ LLVMPositionBuilderAtEnd(ctx->ac.builder, main_fn_body);
real_shader_type = ctx->type;
/* LS is merged into HS (TCS), and ES is merged into GS. */
if (ctx->screen->b.chip_class >= GFX9) {
if (ctx->shader->key.as_ls)
real_shader_type = PIPE_SHADER_TESS_CTRL;
else if (ctx->shader->key.as_es)
real_shader_type = PIPE_SHADER_GEOMETRY;
}
@@ -1403,21 +1387,21 @@ void si_llvm_optimize_module(struct si_shader_context *ctx)
LLVMAddCFGSimplificationPass(gallivm->passmgr);
#if HAVE_LLVM >= 0x0400
/* This is recommended by the instruction combining pass. */
LLVMAddEarlyCSEMemSSAPass(gallivm->passmgr);
#endif
LLVMAddInstructionCombiningPass(gallivm->passmgr);
/* Run the pass */
LLVMRunPassManager(gallivm->passmgr, ctx->gallivm.module);
- LLVMDisposeBuilder(gallivm->builder);
+ LLVMDisposeBuilder(ctx->ac.builder);
LLVMDisposePassManager(gallivm->passmgr);
gallivm_dispose_target_library_info(target_library_info);
}
void si_llvm_dispose(struct si_shader_context *ctx)
{
LLVMDisposeModule(ctx->gallivm.module);
LLVMContextDispose(ctx->gallivm.context);
FREE(ctx->temp_arrays);
ctx->temp_arrays = NULL;
--
2.7.4
More information about the mesa-dev
mailing list