Skip to content
New issue

Have a question about this project? # for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “#”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? # to your account

Plan to support mfma dtype cast? #83

Open
LeiWang1999 opened this issue Jun 19, 2023 · 0 comments
Open

Plan to support mfma dtype cast? #83

LeiWang1999 opened this issue Jun 19, 2023 · 0 comments
Assignees
Labels
hipcc Related to HIPCC

Comments

@LeiWang1999
Copy link

I think float4 and __attribute__((__vector_size__(4 * sizeof(float)))) float are the same thing, but there're no direct conversion of hipcc,

extern "C" __global__ void __launch_bounds__(64) sgemm_16x16x16(float *__restrict__ A, float *__restrict__ B, float *__restrict__ C) {
//   using float4 = __attribute__((__vector_size__(4 * sizeof(float)))) float;
  float Accum[4];
  float MultiA[1];
  float MultiB[1];
  for (int i = 0; i < 4; ++i) {
    Accum[i] = 0.000000e+00f;
  }
  MultiA[0] = A[(((((int)threadIdx.x) & 15) * 4) + (((int)threadIdx.x) >> 4))];
  MultiB[0] = B[(((((int)threadIdx.x) >> 4) * 4) + (((int)threadIdx.x) & 15))];
  *(float4 *)(Accum + 0) =
      (__builtin_amdgcn_mfma_f32_16x16x4f32(
          *((float *)MultiA + 0), *((float *)MultiB + 0),
          *((float4 *)Accum + 0), 0, 0, 0));
  ;
  for (int mma_accum_c_id = 0; mma_accum_c_id < 4; ++mma_accum_c_id) {
    C[((((((int)threadIdx.x) >> 4) * 64) + (mma_accum_c_id * 16)) +
       (((int)threadIdx.x) & 15))] = Accum[mma_accum_c_id];
  }
}

this will throw no viable conversion from 'float4' (aka 'HIP_vector_type<float, 4>') to '__attribute__((__vector_size__(4 * sizeof(float)))) float' (vector of 4 'float' values). This doesn't make sense.

@dgaliffiAMD dgaliffiAMD added the hipcc Related to HIPCC label May 4, 2024
@lamb-j lamb-j transferred this issue from ROCm/HIPCC May 14, 2024
# for free to join this conversation on GitHub. Already have an account? # to comment
Labels
hipcc Related to HIPCC
Projects
None yet
Development

No branches or pull requests

3 participants