Advice for using amd HIP in zig / structure c code for zig c translate

Help
,
1

What I’m currently doing is making me unhappy.

hip.h:

#ifndef HIP
#define HIP

#include <stdint.h>

typedef void *HipStream;

#ifdef __cplusplus
extern "C" {
#endif

uint32_t add(uint32_t a, uint32_t b);

uint32_t createHipStream(HipStream *stream);
void destroyHipStream(HipStream stream);

#ifdef __cplusplus
}
#endif // __cplusplus

#endif // HIP

hip.cpp:

#include "hip.h"

#include <hip/hip_runtime.h>

uint32_t add(uint32_t a, uint32_t b) { return a + b; }

uint32_t createHipStream(HipStream *stream) {
  hipStream_t hip_stream = nullptr;
  const hipError_t error = hipStreamCreate(&hip_stream);
  if (error != hipSuccess) {
    return 1;
  }

  *stream = (void *)hip_stream;

  return 0;
}

void destroyHipStream(HipStream stream) {
  static_cast<void>(hipStreamDestroy((hipStream_t)stream));
}

The cpp code gets compiled into an object file with hipcc, that I link to and I also use the header in translate c:

    const hip_translated_c = b.addTranslateC(.{
        .root_source_file = b.path("libs/hip/hip.h"),
        .optimize = optimize,
        .target = target,
        .link_libc = true,
    });

    const hip_library = b.addLibrary(.{
        .name = "hip",
        .root_module = init: {
            const hip_module = b.createModule(.{
                .optimize = optimize,
                .target = target,
                .link_libcpp = true,
            });

            hip_module.addObjectFile(b.path("libs/hip/hip.o"));

            break :init hip_module;
        },
        .linkage = .static,
    });

    b.installArtifact(hip_library);
    const exe = b.addExecutable(.{
        .name = "ZigWithHip",
        .root_module = init: {
            const exe_module = b.createModule(.{
                .root_source_file = b.path("src/main.zig"),
                .target = target,
                .optimize = optimize,
                .imports = &.{
                    .{ .name = "hip", .module = hip_translated_c.createModule() },
                },
            });

            exe_module.linkLibrary(hip_library);
            exe_module.linkSystemLibrary("amdhip64", .{});

            break :init exe_module;
        },
    });

    b.installArtifact(exe);


// ETC..

Any tips on how to make the cpp side better so it translates “nicely” to zig? I tried to avoid having any HIP stuff in the outfacing headers by casting to void ptrs, but should I should I just include the HIP header in the hip.h so they also get translated? Any examples or links are greatly appreciated.

Thank you!

2

A library exposing just a void* context pointer is a common pattern in C/C++.

What part of this specifically is it that you find cumbersome or “bad”?

You don’t show any Zig code consuming the library. Perhaps it would help to show from the Zig side whatever it is that you’re not liking.

2 Likes
3

Well it gets translated in the build phase, so what you see in the header is what you get. I was just a little unsure about if there was away to make the translation nicer. My post was maybe a little vague. It should be clearer now, hopefully.

4

I’m still not clear. Are you asking about how the translation and linking is set up in your build.zig?

5

No, more of how to make it more C code more convenient to use in zig. I know you can’t do much since the translation is a translation, but is there pitfalls or helpful patterns etc. To make using the C code a little nicer to use. Or do I just have to add a zig wrapper to the wrapper. It’s just kind of annoying to have two layers of wrappers if that is the case.

6

For example errors, what do you usually do here:

    const HipError = enum(u32) {
        ok = 0,
        _,

        pub inline fn toError(value: u32) @This() {
            return @enumFromInt(value);
        }
    };

    var hipStream: hip.HipStream = undefined;
    if (HipError.toError(hip.createHipStream(&hipStream)) != .ok) return error.CreateHipStreamFailed;
    defer hip.destroyHipStream(hipStream);

    var hipStream: hip.HipStream = undefined;
    if (hip.createHipStream(&hipStream) != 0) return error.CreateHipStreamFailed;
    defer hip.destroyHipStream(hipStream);

    const hip_ok: u32 = 0;
    var hipStream: hip.HipStream = undefined;
    if (hip.createHipStream(&hipStream) != hip_ok) return error.CreateHipStreamFailed;
    defer hip.destroyHipStream(hipStream);

Everything is making me equally unhappy. They do all work, and I’m kind of leaning into the enum but having to wrap it with toError() is sub optimal.

EDIT, maybe?:

pub fn bubbleWrap(value: u32) !void {
    switch (value) {
        0 => return,
        1 => return error.CreateHipStreamFailed,
        _ => return error.UnknownHipFail,
    }
}

var hipStream: hip.HipStream = undefined;
try bubbleWrap(hip.createHipStream(&hipStream)));
7

I’m probably just going to wrap this with a normal function call, but any other ideas on how to call kernels?

// Example pseudo code introducing hipLaunchKernelGGL:
__global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N)
{
...
}

MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a,b,c,n);
// Alternatively, kernel can be launched by
// hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n);

From: HIP docs