Skip to content

support __hip_bfloat16 type alias (fixes #65)#113

Open
GauthamMK-0 wants to merge 1 commit into
Zaneham:masterfrom
GauthamMK-0:master
Open

support __hip_bfloat16 type alias (fixes #65)#113
GauthamMK-0 wants to merge 1 commit into
Zaneham:masterfrom
GauthamMK-0:master

Conversation

@GauthamMK-0

Copy link
Copy Markdown

Fix: HIP __hip_bfloat16 Type Alias Support (Issue #65)

Problem

Previously, BarraCUDA only recognized the CUDA native __nv_bfloat16 type. When encountering HIP's __hip_bfloat16 type alias, it was passed as int

Solution

Added __hip_bfloat16 type alias resolution mappings within the semantic analyzer and intermediate representation

  • src/fe/sema.c
  • src/ir/bir_lower.c

Verification

Test Code Sample

The following snippet was used to verify semantic type inference of __hip_bfloat16

#include <hip/hip_runtime.h>
#include <hip/hip_bfloat16.h>

__global__ void test_bf16(__hip_bfloat16 *out, __hip_bfloat16 a, __hip_bfloat16 b) {
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    out[id] = a + b;
}

Before __hip_bfloat16 mapping:

; Sema: 19 types interned, 1 symbols, 0 structs, 0 error(s)
; Symbols:
;   test_bf16            func       func(int*, int, int) -> void
; Annotated expressions:
;   node[20] var_decl @ 5:12 → int
;   node[22] member @ 5:22 → unsigned int
;   node[24] binary @ 5:25 → unsigned int
;   node[26] member @ 5:35 → unsigned int
;   node[28] binary @ 5:38 → unsigned int
;   node[30] member @ 5:49 → unsigned int
;   node[32] ident @ 6:5 → int*
;   node[33] subscript @ 6:8 → int
;   node[34] ident @ 6:9 → int
;   node[35] binary @ 6:13 → int
;   node[36] ident @ 6:15 → int
;   node[37] binary @ 6:17 → int
;   node[38] ident @ 6:19 → int
; 13 nodes annotated

After __hip_bfloat16 mapping:

; Sema: 20 types interned, 1 symbols, 0 structs, 0 error(s)
; Symbols:
;   test_bf16            func       func(__bfloat16*, __bfloat16, __bfloat16) -> void
; Annotated expressions:
;   node[20] var_decl @ 5:12 → int
;   node[22] member @ 5:22 → unsigned int
;   node[24] binary @ 5:25 → unsigned int
;   node[26] member @ 5:35 → unsigned int
;   node[28] binary @ 5:38 → unsigned int
;   node[30] member @ 5:49 → unsigned int
;   node[32] ident @ 6:5 → __bfloat16*
;   node[33] subscript @ 6:8 → __bfloat16
;   node[34] ident @ 6:9 → int
;   node[35] binary @ 6:13 → __bfloat16
;   node[36] ident @ 6:15 → __bfloat16
;   node[37] binary @ 6:17 → __bfloat16
;   node[38] ident @ 6:19 → __bfloat16
; 13 nodes annotated

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant