diff options
Diffstat (limited to 'src/codegen')
| -rw-r--r-- | src/codegen/c.zig | 55 | ||||
| -rw-r--r-- | src/codegen/llvm.zig | 86 | ||||
| -rw-r--r-- | src/codegen/llvm/bindings.zig | 3 |
3 files changed, 117 insertions, 27 deletions
diff --git a/src/codegen/c.zig b/src/codegen/c.zig index 66b0ef9e75..6577089806 100644 --- a/src/codegen/c.zig +++ b/src/codegen/c.zig @@ -1071,7 +1071,7 @@ pub const DeclGen = struct { const extern_fn = val.castTag(.extern_fn).?.data; try dg.renderDeclName(writer, extern_fn.owner_decl, 0); }, - .int_u64, .one => { + .int_u64, .one, .int_big_positive, .lazy_align, .lazy_size => { try writer.writeAll("(("); try dg.renderType(writer, ty); return writer.print("){x})", .{try dg.fmtIntLiteral(Type.usize, val, .Other)}); @@ -2997,6 +2997,11 @@ fn genBodyInner(f: *Function, body: []const Air.Inst.Index) error{ AnalysisFail, .c_va_arg => try airCVaArg(f, inst), .c_va_end => try airCVaEnd(f, inst), .c_va_copy => try airCVaCopy(f, inst), + + .work_item_id, + .work_group_size, + .work_group_id, + => unreachable, // zig fmt: on }; if (result_value == .new_local) { @@ -3594,10 +3599,6 @@ fn airStore(f: *Function, inst: Air.Inst.Index) !CValue { const ptr_ty = f.air.typeOf(bin_op.lhs); const ptr_scalar_ty = ptr_ty.scalarType(); const ptr_info = ptr_scalar_ty.ptrInfo().data; - if (!ptr_info.pointee_type.hasRuntimeBitsIgnoreComptime()) { - try reap(f, inst, &.{ bin_op.lhs, bin_op.rhs }); - return .none; - } const ptr_val = try f.resolveInst(bin_op.lhs); const src_ty = f.air.typeOf(bin_op.rhs); @@ -4458,9 +4459,7 @@ fn airBr(f: *Function, inst: Air.Inst.Index) !CValue { fn airBitcast(f: *Function, inst: Air.Inst.Index) !CValue { const ty_op = f.air.instructions.items(.data)[inst].ty_op; const dest_ty = f.air.typeOfIndex(inst); - // No IgnoreComptime until Sema stops giving us garbage Air. - // https://github.com/ziglang/zig/issues/13410 - if (f.liveness.isUnused(inst) or !dest_ty.hasRuntimeBits()) { + if (f.liveness.isUnused(inst)) { try reap(f, inst, &.{ty_op.operand}); return .none; } @@ -6854,17 +6853,35 @@ fn airAggregateInit(f: *Function, inst: Air.Inst.Index) !CValue { switch (inst_ty.zigTypeTag()) { .Array, .Vector => { const elem_ty = inst_ty.childType(); - for (resolved_elements, 0..) |element, i| { - try f.writeCValue(writer, local, .Other); - try writer.print("[{d}] = ", .{i}); - try f.writeCValue(writer, element, .Other); - try writer.writeAll(";\n"); - } - if (inst_ty.sentinel()) |sentinel| { - try f.writeCValue(writer, local, .Other); - try writer.print("[{d}] = ", .{resolved_elements.len}); - try f.object.dg.renderValue(writer, elem_ty, sentinel, .Other); - try writer.writeAll(";\n"); + + const is_array = lowersToArray(elem_ty, target); + const need_memcpy = is_array; + if (need_memcpy) { + for (resolved_elements, 0..) |element, i| { + try writer.writeAll("memcpy("); + try f.writeCValue(writer, local, .Other); + try writer.print("[{d}]", .{i}); + try writer.writeAll(", "); + try f.writeCValue(writer, element, .Other); + try writer.writeAll(", sizeof("); + try f.renderType(writer, elem_ty); + try writer.writeAll("))"); + try writer.writeAll(";\n"); + } + assert(inst_ty.sentinel() == null); + } else { + for (resolved_elements, 0..) |element, i| { + try f.writeCValue(writer, local, .Other); + try writer.print("[{d}] = ", .{i}); + try f.writeCValue(writer, element, .Other); + try writer.writeAll(";\n"); + } + if (inst_ty.sentinel()) |sentinel| { + try f.writeCValue(writer, local, .Other); + try writer.print("[{d}] = ", .{resolved_elements.len}); + try f.object.dg.renderValue(writer, elem_ty, sentinel, .Other); + try writer.writeAll(";\n"); + } } }, .Struct => switch (inst_ty.containerLayout()) { diff --git a/src/codegen/llvm.zig b/src/codegen/llvm.zig index 18237efe65..233ec21ac1 100644 --- a/src/codegen/llvm.zig +++ b/src/codegen/llvm.zig @@ -2724,6 +2724,9 @@ pub const DeclGen = struct { if (comp.bin_file.options.llvm_cpu_features) |s| { llvm_fn.addFunctionAttr("target-features", s); } + if (comp.getTarget().cpu.arch.isBpf()) { + llvm_fn.addFunctionAttr("no-builtins", ""); + } } fn resolveGlobalDecl(dg: *DeclGen, decl_index: Module.Decl.Index) Error!*llvm.Value { @@ -3402,7 +3405,7 @@ pub const DeclGen = struct { }; return dg.context.constStruct(&fields, fields.len, .False); }, - .int_u64, .one, .int_big_positive => { + .int_u64, .one, .int_big_positive, .lazy_align, .lazy_size => { const llvm_usize = try dg.lowerType(Type.usize); const llvm_int = llvm_usize.constInt(tv.val.toUnsignedInt(target), .False); return llvm_int.constIntToPtr(try dg.lowerType(tv.ty)); @@ -3820,6 +3823,8 @@ pub const DeclGen = struct { const field_ty = union_obj.fields.values()[field_index].ty; if (union_obj.layout == .Packed) { + if (!field_ty.hasRuntimeBits()) + return llvm_union_ty.constNull(); const non_int_val = try lowerValue(dg, .{ .ty = field_ty, .val = tag_and_val.val }); const ty_bit_size = @intCast(u16, field_ty.bitSize(target)); const small_int_ty = dg.context.intType(ty_bit_size); @@ -4745,6 +4750,10 @@ pub const FuncGen = struct { .c_va_copy => try self.airCVaCopy(inst), .c_va_end => try self.airCVaEnd(inst), .c_va_start => try self.airCVaStart(inst), + + .work_item_id => try self.airWorkItemId(inst), + .work_group_size => try self.airWorkGroupSize(inst), + .work_group_id => try self.airWorkGroupId(inst), // zig fmt: on }; if (opt_value) |val| { @@ -5561,7 +5570,7 @@ pub const FuncGen = struct { return fg.loadByRef(payload_ptr, payload_ty, payload_ty.abiAlignment(target), false); } - const load_inst = fg.builder.buildLoad(payload_ptr.getGEPResultElementType(), payload_ptr, ""); + const load_inst = fg.builder.buildLoad(err_union_llvm_ty.structGetTypeAtIndex(offset), payload_ptr, ""); load_inst.setAlignment(payload_ty.abiAlignment(target)); return load_inst; } @@ -6795,7 +6804,7 @@ pub const FuncGen = struct { return self.loadByRef(payload_ptr, payload_ty, payload_ty.abiAlignment(target), false); } - const load_inst = self.builder.buildLoad(payload_ptr.getGEPResultElementType(), payload_ptr, ""); + const load_inst = self.builder.buildLoad(err_union_llvm_ty.structGetTypeAtIndex(offset), payload_ptr, ""); load_inst.setAlignment(payload_ty.abiAlignment(target)); return load_inst; } @@ -8212,7 +8221,6 @@ pub const FuncGen = struct { const dest_ptr = try self.resolveInst(bin_op.lhs); const ptr_ty = self.air.typeOf(bin_op.lhs); const operand_ty = ptr_ty.childType(); - if (!operand_ty.isFnOrHasRuntimeBitsIgnoreComptime()) return null; // TODO Sema should emit a different instruction when the store should // possibly do the safety 0xaa bytes for undefined. @@ -8583,7 +8591,7 @@ pub const FuncGen = struct { } const tag_index = @boolToInt(layout.tag_align < layout.payload_align); const tag_field_ptr = self.builder.buildStructGEP(llvm_un_ty, union_handle, tag_index, ""); - return self.builder.buildLoad(tag_field_ptr.getGEPResultElementType(), tag_field_ptr, ""); + return self.builder.buildLoad(llvm_un_ty.structGetTypeAtIndex(tag_index), tag_field_ptr, ""); } else { if (layout.payload_size == 0) { return union_handle; @@ -9567,6 +9575,74 @@ pub const FuncGen = struct { return self.builder.buildAddrSpaceCast(operand, llvm_dest_ty, ""); } + fn amdgcnWorkIntrinsic(self: *FuncGen, dimension: u32, default: u32, comptime basename: []const u8) !?*llvm.Value { + const llvm_u32 = self.context.intType(32); + + const llvm_fn_name = switch (dimension) { + 0 => basename ++ ".x", + 1 => basename ++ ".y", + 2 => basename ++ ".z", + else => return llvm_u32.constInt(default, .False), + }; + + const args: [0]*llvm.Value = .{}; + const llvm_fn = self.getIntrinsic(llvm_fn_name, &.{}); + return self.builder.buildCall(llvm_fn.globalGetValueType(), llvm_fn, &args, args.len, .Fast, .Auto, ""); + } + + fn airWorkItemId(self: *FuncGen, inst: Air.Inst.Index) !?*llvm.Value { + if (self.liveness.isUnused(inst)) return null; + + const target = self.dg.module.getTarget(); + assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures + + const pl_op = self.air.instructions.items(.data)[inst].pl_op; + const dimension = pl_op.payload; + return self.amdgcnWorkIntrinsic(dimension, 0, "llvm.amdgcn.workitem.id"); + } + + fn airWorkGroupSize(self: *FuncGen, inst: Air.Inst.Index) !?*llvm.Value { + if (self.liveness.isUnused(inst)) return null; + + const target = self.dg.module.getTarget(); + assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures + + const pl_op = self.air.instructions.items(.data)[inst].pl_op; + const dimension = pl_op.payload; + const llvm_u32 = self.context.intType(32); + if (dimension >= 3) { + return llvm_u32.constInt(1, .False); + } + + // Fetch the dispatch pointer, which points to this structure: + // https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/adae6c61e10d371f7cbc3d0e94ae2c070cab18a4/src/inc/hsa.h#L2913 + const llvm_fn = self.getIntrinsic("llvm.amdgcn.dispatch.ptr", &.{}); + const args: [0]*llvm.Value = .{}; + const dispatch_ptr = self.builder.buildCall(llvm_fn.globalGetValueType(), llvm_fn, &args, args.len, .Fast, .Auto, ""); + dispatch_ptr.setAlignment(4); + + // Load the work_group_* member from the struct as u16. + // Just treat the dispatch pointer as an array of u16 to keep things simple. + const offset = 2 + dimension; + const index = [_]*llvm.Value{llvm_u32.constInt(offset, .False)}; + const llvm_u16 = self.context.intType(16); + const workgroup_size_ptr = self.builder.buildInBoundsGEP(llvm_u16, dispatch_ptr, &index, index.len, ""); + const workgroup_size = self.builder.buildLoad(llvm_u16, workgroup_size_ptr, ""); + workgroup_size.setAlignment(2); + return workgroup_size; + } + + fn airWorkGroupId(self: *FuncGen, inst: Air.Inst.Index) !?*llvm.Value { + if (self.liveness.isUnused(inst)) return null; + + const target = self.dg.module.getTarget(); + assert(target.cpu.arch == .amdgcn); // TODO is to port this function to other GPU architectures + + const pl_op = self.air.instructions.items(.data)[inst].pl_op; + const dimension = pl_op.payload; + return self.amdgcnWorkIntrinsic(dimension, 0, "llvm.amdgcn.workgroup.id"); + } + fn getErrorNameTable(self: *FuncGen) !*llvm.Value { if (self.dg.object.error_name_table) |table| { return table; diff --git a/src/codegen/llvm/bindings.zig b/src/codegen/llvm/bindings.zig index 4286be5e65..c5a6d81ff3 100644 --- a/src/codegen/llvm/bindings.zig +++ b/src/codegen/llvm/bindings.zig @@ -254,9 +254,6 @@ pub const Value = opaque { pub const addFunctionAttr = ZigLLVMAddFunctionAttr; extern fn ZigLLVMAddFunctionAttr(Fn: *Value, attr_name: [*:0]const u8, attr_value: [*:0]const u8) void; - pub const getGEPResultElementType = ZigLLVMGetGEPResultElementType; - extern fn ZigLLVMGetGEPResultElementType(GEP: *Value) *Type; - pub const addByValAttr = ZigLLVMAddByValAttr; extern fn ZigLLVMAddByValAttr(Fn: *Value, ArgNo: c_uint, type: *Type) void; }; |