Skip to content

Implement hipSetKernelAttributeDevice API #10542

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
vstojilj merged 2 commits into from
Oct 5, 2026
Merged

Implement hipSetKernelAttributeDevice API #10542

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
5186ebd
Implement hipSetKernelAttributeDevice API
vstojilj Aug 21, 2026
ce8b8c6
fix(hip): copy in-memory library images for secondary-device loads
vstojilj Sep 11, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 7 additions & 2 deletions projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,7 @@
#define HIP_API_TABLE_STEP_VERSION 0
#define HIP_COMPILER_API_TABLE_STEP_VERSION 0
#define HIP_TOOLS_API_TABLE_STEP_VERSION 1
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 35
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 36

// HIP API interface
// HIP compiler dispatch functions
Expand Down Expand Up @@ -1189,6 +1189,8 @@ typedef hipError_t (*t_hipModuleEnumerateFunctions)(hipFunction_t* functions,
// rocprofiler-sdk/hip/api_args.h), and these enums are declared without a typedef.
typedef hipError_t (*t_hipDeviceFlushGPUDirectRDMAWrites)(
enum hipFlushGPUDirectRDMAWritesTarget target, enum hipFlushGPUDirectRDMAWritesScope scope);
typedef hipError_t (*t_hipKernelSetAttributeForDevice)(hipKernel_t kernel, hipFuncAttribute attr,
int value, int device);
// HIP Compiler dispatch table
struct HipCompilerDispatchTable {
// HIP_COMPILER_API_TABLE_STEP_VERSION == 0
Expand Down Expand Up @@ -1848,8 +1850,11 @@ struct HipDispatchTable {
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 35
t_hipDeviceFlushGPUDirectRDMAWrites hipDeviceFlushGPUDirectRDMAWrites_fn;

// DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 36
t_hipKernelSetAttributeForDevice hipKernelSetAttributeForDevice_fn;

// DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 37

// ******************************************************************************************* //
//
Expand Down
29 changes: 28 additions & 1 deletion projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -508,7 +508,8 @@ enum hip_api_id_t {
HIP_API_ID_hipInitDevice = 486,
HIP_API_ID_hipModuleEnumerateFunctions = 487,
HIP_API_ID_hipDeviceFlushGPUDirectRDMAWrites = 488,
HIP_API_ID_LAST = 488,
HIP_API_ID_hipKernelSetAttributeForDevice = 489,
HIP_API_ID_LAST = 489,


HIP_API_ID_hipBindTexture = HIP_API_ID_NONE,
Expand Down Expand Up @@ -795,6 +796,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipKernelGetName: return "hipKernelGetName";
case HIP_API_ID_hipKernelGetParamInfo: return "hipKernelGetParamInfo";
case HIP_API_ID_hipKernelSetAttribute: return "hipKernelSetAttribute";
case HIP_API_ID_hipKernelSetAttributeForDevice: return "hipKernelSetAttributeForDevice";
case HIP_API_ID_hipLaunchByPtr: return "hipLaunchByPtr";
case HIP_API_ID_hipLaunchCooperativeKernel: return "hipLaunchCooperativeKernel";
case HIP_API_ID_hipLaunchCooperativeKernelMultiDevice: return "hipLaunchCooperativeKernelMultiDevice";
Expand Down Expand Up @@ -1277,6 +1279,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipKernelGetName", name) == 0) return HIP_API_ID_hipKernelGetName;
if (strcmp("hipKernelGetParamInfo", name) == 0) return HIP_API_ID_hipKernelGetParamInfo;
if (strcmp("hipKernelSetAttribute", name) == 0) return HIP_API_ID_hipKernelSetAttribute;
if (strcmp("hipKernelSetAttributeForDevice", name) == 0) return HIP_API_ID_hipKernelSetAttributeForDevice;
if (strcmp("hipLaunchByPtr", name) == 0) return HIP_API_ID_hipLaunchByPtr;
if (strcmp("hipLaunchCooperativeKernel", name) == 0) return HIP_API_ID_hipLaunchCooperativeKernel;
if (strcmp("hipLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipLaunchCooperativeKernelMultiDevice;
Expand Down Expand Up @@ -2964,6 +2967,12 @@ typedef struct hip_api_data_s {
hipKernel_t kernel;
hipDevice_t dev;
} hipKernelSetAttribute;
struct {
hipKernel_t kernel;
hipFuncAttribute attr;
int value;
int device;
} hipKernelSetAttributeForDevice;
struct {
const void* hostFunction;
} hipLaunchByPtr;
Expand Down Expand Up @@ -5892,6 +5901,13 @@ typedef struct hip_api_data_s {
cb_data.args.hipKernelSetAttribute.kernel = (hipKernel_t)kernel; \
cb_data.args.hipKernelSetAttribute.dev = (hipDevice_t)dev; \
};
// hipKernelSetAttributeForDevice[('hipKernel_t', 'kernel'), ('hipFuncAttribute', 'attr'), ('int', 'value'), ('int', 'device')]
#define INIT_hipKernelSetAttributeForDevice_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipKernelSetAttributeForDevice.kernel = (hipKernel_t)kernel; \
cb_data.args.hipKernelSetAttributeForDevice.attr = (hipFuncAttribute)attr; \
cb_data.args.hipKernelSetAttributeForDevice.value = (int)value; \
cb_data.args.hipKernelSetAttributeForDevice.device = (int)device; \
};
// hipLaunchByPtr[('const void*', 'hostFunction')]
#define INIT_hipLaunchByPtr_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipLaunchByPtr.hostFunction = (const void*)hostFunction; \
Expand Down Expand Up @@ -8443,6 +8459,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipKernelSetAttribute[('hipFunction_attribute', 'attrib'), ('int', 'value'), ('hipKernel_t', 'kernel'), ('hipDevice_t', 'dev')]
case HIP_API_ID_hipKernelSetAttribute:
break;
// hipKernelSetAttributeForDevice[('hipKernel_t', 'kernel'), ('hipFuncAttribute', 'attr'), ('int', 'value'), ('int', 'device')]
case HIP_API_ID_hipKernelSetAttributeForDevice:
break;
// hipLaunchByPtr[('const void*', 'hostFunction')]
case HIP_API_ID_hipLaunchByPtr:
break;
Expand Down Expand Up @@ -11313,6 +11332,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", dev="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelSetAttribute.dev);
oss << ")";
break;
case HIP_API_ID_hipKernelSetAttributeForDevice:
oss << "hipKernelSetAttributeForDevice(";
oss << "kernel="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelSetAttributeForDevice.kernel);
oss << ", attr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelSetAttributeForDevice.attr);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelSetAttributeForDevice.value);
oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipKernelSetAttributeForDevice.device);
oss << ")";
break;
case HIP_API_ID_hipLaunchByPtr:
oss << "hipLaunchByPtr(";
oss << "hostFunction="; roctracer::hip_support::detail::operator<<(oss, data->args.hipLaunchByPtr.hostFunction);
Expand Down
1 change: 1 addition & 0 deletions projects/clr/hipamd/src/amdhip.def.in
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -543,6 +543,7 @@ hipMemGetMemPool
hipMipmappedArrayGetMemoryRequirements
hipKernelGetAttribute
hipKernelSetAttribute
hipKernelSetAttributeForDevice
hipKernelGetFunction
hipOccupancyMaxPotentialClusterSize
hipOccupancyMaxActiveClusters
Expand Down
11 changes: 8 additions & 3 deletions projects/clr/hipamd/src/hip_api_trace.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -940,9 +940,11 @@ hipError_t hipExecutionCtxWaitEvent(hipExecutionCtx_t ctx, hipEvent_t event);
hipError_t hipMemGetDefaultMemPool(hipMemPool_t* memPool, hipMemLocation* location,
hipMemAllocationType type);
hipError_t hipModuleEnumerateFunctions(hipFunction_t* functions, unsigned int numFunctions,
hipModule_t mod);
hipModule_t mod);
hipError_t hipDeviceFlushGPUDirectRDMAWrites(hipFlushGPUDirectRDMAWritesTarget target,
hipFlushGPUDirectRDMAWritesScope scope);
hipError_t hipKernelSetAttributeForDevice(hipKernel_t kernel, hipFuncAttribute attr, int value,
int device);
} // namespace hip

namespace hip {
Expand Down Expand Up @@ -1511,6 +1513,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipMemGetMemPool_fn = hip::hipMemGetMemPool;
ptrDispatchTable->hipKernelGetAttribute_fn = hip::hipKernelGetAttribute;
ptrDispatchTable->hipKernelSetAttribute_fn = hip::hipKernelSetAttribute;
ptrDispatchTable->hipKernelSetAttributeForDevice_fn = hip::hipKernelSetAttributeForDevice;
ptrDispatchTable->hipKernelGetFunction_fn = hip::hipKernelGetFunction;
ptrDispatchTable->hipGreenCtxCreate_fn = hip::hipGreenCtxCreate;
ptrDispatchTable->hipExecutionCtxDestroy_fn = hip::hipExecutionCtxDestroy;
Expand Down Expand Up @@ -2274,15 +2277,17 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipInitDevice_fn, 543);
HIP_ENFORCE_ABI(HipDispatchTable, hipModuleEnumerateFunctions_fn, 544);
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 35
HIP_ENFORCE_ABI(HipDispatchTable, hipDeviceFlushGPUDirectRDMAWrites_fn, 545);
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 36
HIP_ENFORCE_ABI(HipDispatchTable, hipKernelSetAttributeForDevice_fn, 546);
// if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below
// will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.:
//
// HIP_ENFORCE_ABI(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 9) <- 8 + 1 = 9
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 546)
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 547)

static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 35,
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 36,
"If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
"pointers and then update this check so it is true");
#endif
27 changes: 15 additions & 12 deletions projects/clr/hipamd/src/hip_code_object.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,21 +24,24 @@ hipError_t ihipFree(void* ptr);
// forward declaration of methods required for managed variables
hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0);

hipError_t DynCO::loadCodeObject(const char* fname, const void* image) {
hipError_t DynCO::loadCodeObject(const char* fname, const void* image, bool init_global_vars,
std::vector<char>* image_storage) {
std::scoped_lock lock(dclock_);

// Number of devices = 1 in dynamic code object
fb_info_ = new FatBinaryInfo(fname, image);
std::vector<hip::Device*> devices = {g_devices[ihipGetDevice()]};
IHIP_RETURN_ONFAIL(fb_info_->ExtractFatBinaryUsingCOMGR(devices));
std::vector<hip::Device*> devices = {g_devices[device_id_]};
IHIP_RETURN_ONFAIL(fb_info_->ExtractFatBinaryUsingCOMGR(devices, image_storage));

// No Lazy loading for DynCO
IHIP_RETURN_ONFAIL(fb_info_->BuildProgram(ihipGetDevice()));
IHIP_RETURN_ONFAIL(fb_info_->BuildProgram(device_id_));

module_ = fb_info_->Module(device_id_);

// Define Global variables
IHIP_RETURN_ONFAIL(populateDynGlobalVars());
if (init_global_vars) {
// Primary library/module loads own global and managed-variable state.
IHIP_RETURN_ONFAIL(populateDynGlobalVars());
}

// Define Global functions
IHIP_RETURN_ONFAIL(populateDynGlobalFuncs());
Expand Down Expand Up @@ -144,7 +147,7 @@ hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, const std::string& func_name)
}

/* See if this could be solved */
return it->second->GetDynFunc(hfunc, module_);
return it->second->GetDynFunc(hfunc, module_, device_id_);
}

hipError_t DynCO::getFuncCount(unsigned int* count) {
Expand All @@ -166,7 +169,7 @@ hipError_t DynCO::enumerateFunctions(hipFunction_t* functions, unsigned int numF
break;
}
hipFunction_t hfunc = nullptr;
auto ret = kv.second->GetDynFunc(&hfunc, module_);
auto ret = kv.second->GetDynFunc(&hfunc, module_, device_id_);
if (ret != hipSuccess) {
return ret;
}
Expand Down Expand Up @@ -234,8 +237,8 @@ hipError_t DynCO::populateDynGlobalVars() {
std::vector<std::string> var_names;
std::string managedVarExt = ".managed";
// For Dynamic Modules there is only one hipFatBinaryDevInfo_
device::Program* dev_program = fb_info_->GetProgram(ihipGetDevice())
->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
device::Program* dev_program =
fb_info_->GetProgram(device_id_)->getDeviceProgram(*g_devices[device_id_]->devices()[0]);

if (!dev_program->getGlobalVarFromCodeObj(&var_names)) {
LogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x", module_);
Expand All @@ -261,8 +264,8 @@ hipError_t DynCO::populateDynGlobalFuncs() {
std::scoped_lock lock(dclock_);

std::vector<std::string> func_names;
device::Program* dev_program = fb_info_->GetProgram(ihipGetDevice())
->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
device::Program* dev_program =
fb_info_->GetProgram(device_id_)->getDeviceProgram(*g_devices[device_id_]->devices()[0]);

// Get all the global func names from COMGR
if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) {
Expand Down
11 changes: 8 additions & 3 deletions projects/clr/hipamd/src/hip_code_object.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,11 +94,16 @@ class DynCO : public CodeObject {
std::recursive_mutex dclock_;

public:
DynCO() : device_id_(ihipGetDevice()), fb_info_(nullptr), module_(nullptr) {}
explicit DynCO(int device_id = ihipGetDevice())
: device_id_(device_id), fb_info_(nullptr), module_(nullptr) {}
virtual ~DynCO();

// LoadsCodeObject and its data
hipError_t loadCodeObject(const char* fname, const void* image = nullptr);
// Primary device loads code object and initialize global and managed variables. In case we need
// to set attribute for a device that's not current, we need to load the code object for it, but
// don't need to initialize global/managed-variable state.
hipError_t loadCodeObject(const char* fname, const void* image = nullptr,
bool init_global_vars = true,
std::vector<char>* image_storage = nullptr);
hipModule_t getModule() const { return module_; };

// Device the code object was loaded for at construction. Callers that key
Expand Down
Loading
Loading