From a19f98a134d74ab3d8d9ee9cfa61f748b3adc69d Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sat, 26 Aug 2023 15:48:03 +0200 Subject: [PATCH] ac: drop 64 bit handling for cl workgroup intrinsics MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Reviewed-by: Marek Olšák Part-of: --- src/amd/llvm/ac_nir_to_llvm.c | 4 ---- src/amd/vulkan/meta/radv_meta_buffer.c | 12 ++++++------ 2 files changed, 6 insertions(+), 10 deletions(-) diff --git a/src/amd/llvm/ac_nir_to_llvm.c b/src/amd/llvm/ac_nir_to_llvm.c index 3d62d8da4ef..d857d4a85d3 100644 --- a/src/amd/llvm/ac_nir_to_llvm.c +++ b/src/amd/llvm/ac_nir_to_llvm.c @@ -3057,8 +3057,6 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins values[i] = ctx->args->workgroup_ids[i].used ? ac_get_arg(&ctx->ac, ctx->args->workgroup_ids[i]) : ctx->ac.i32_0; - if (instr->def.bit_size == 64) - values[i] = LLVMBuildZExt(ctx->ac.builder, values[i], ctx->ac.i64, ""); } result = ac_build_gather_values(&ctx->ac, values, 3); @@ -3156,8 +3154,6 @@ static bool visit_intrinsic(struct ac_nir_context *ctx, nir_intrinsic_instr *ins result = ac_build_load_invariant(&ctx->ac, ac_get_ptr_arg(&ctx->ac, ctx->args, ctx->args->num_work_groups), ctx->ac.i32_0); } - if (instr->def.bit_size == 64) - result = LLVMBuildZExt(ctx->ac.builder, result, LLVMVectorType(ctx->ac.i64, 3), ""); break; case nir_intrinsic_load_local_invocation_index: result = visit_load_local_invocation_index(ctx); diff --git a/src/amd/vulkan/meta/radv_meta_buffer.c b/src/amd/vulkan/meta/radv_meta_buffer.c index 4f95749d27a..76f7d4e5111 100644 --- a/src/amd/vulkan/meta/radv_meta_buffer.c +++ b/src/amd/vulkan/meta/radv_meta_buffer.c @@ -15,9 +15,9 @@ build_buffer_fill_shader(struct radv_device *dev) nir_def *max_offset = nir_channel(&b, pconst, 2); nir_def *data = nir_swizzle(&b, nir_channel(&b, pconst, 3), (unsigned[]){0, 0, 0, 0}, 4); - nir_def *global_id = nir_iadd( - &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), - nir_load_local_invocation_index(&b)); + nir_def *global_id = + nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), + nir_load_local_invocation_index(&b)); nir_def *offset = nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset); nir_def *dst_addr = nir_iadd(&b, buffer_addr, nir_u2u64(&b, offset)); @@ -37,9 +37,9 @@ build_buffer_copy_shader(struct radv_device *dev) nir_def *src_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b0011)); nir_def *dst_addr = nir_pack_64_2x32(&b, nir_channels(&b, pconst, 0b1100)); - nir_def *global_id = nir_iadd( - &b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), - nir_load_local_invocation_index(&b)); + nir_def *global_id = + nir_iadd(&b, nir_imul_imm(&b, nir_channel(&b, nir_load_workgroup_id(&b), 0), b.shader->info.workgroup_size[0]), + nir_load_local_invocation_index(&b)); nir_def *offset = nir_u2u64(&b, nir_imin(&b, nir_imul_imm(&b, global_id, 16), max_offset));