ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 17:19:12 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#9333 (commit 78aa8c7)

Browse Source 
feat(ck): Add gfx90c build support

## Motivation

Needed for https://github.com/ROCm/TheRock/issues/3818.

## Technical Details

When building torch w/ TheRock CI (`build_prod_wheels.py`) on gfx90c,
was running into the generic error below compiling
`amd_buffer_addressing_builtins.hpp`
```
  composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:33:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
  ...:61:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
  ...:73:22: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'
  4 errors generated when compiling for gfx90c.
  ```
 Adding `__gfx90c__` to the arch definition list resolves this.

## Test Plan

Build torch (`build_prod_wheels.py`) for gfx90c before and after patch.

## Test Result

- Before: Fails with `amd_buffer_addressing_builtins.hpp:33:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'`

- After: Build passes

## Submission Checklist

- [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This commit is contained in:
Harkirat Gill
2026-07-16 19:04:04 +00:00
committed by assistant-librarian[bot]
parent 9d8386b920
commit 973a0bfd81
2 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
include/ck/ck.hpp
View File

@@ -85,7 +85,8 @@
// buffer resource
#ifndef __HIP_DEVICE_COMPILE__ // for host code
#define CK_BUFFER_RESOURCE_3RD_DWORD -1
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx906__) || defined(__gfx9__)
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx90c__) || defined(__gfx906__) || \
defined(__gfx9__)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
#elif defined(__gfx101__) || defined(__gfx103__)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
@@ -102,7 +103,7 @@
#define CK_USE_AMD_V_FMAC_F32
#define CK_USE_AMD_V_DOT2_F32_F16
#define CK_USE_AMD_V_DOT4_I32_I8
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx101__)
#elif defined(__gfx803__) || defined(__gfx900__) || defined(__gfx90c__) || defined(__gfx101__)
#define CK_USE_AMD_V_MAC_F32
#elif defined(__gfx11__) || defined(__gfx120__)
#define CK_USE_AMD_V_FMAC_F32

1
test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data.cpp
View File

@@ -146,7 +146,6 @@ TYPED_TEST(TestGroupedConvndBwdData2d, Test2D)
{2, 1, 4, 128, 64, {3, 3}, {28, 28}, {3, 3}, {1, 1}, {1, 1}, {1, 1}});
this->conv_params.push_back(
{2, 1, 2, 64, 64, {6, 6}, {28, 28}, {6, 6}, {1, 1}, {1, 1}, {1, 1}});
this->template Run<2>();
}