0
11
0
0
DTK-24.04.1版本动态并行
2024-06-14 12:19:55
60
请问下，DTK-24.04.1版本的动态并行的功能如何使用呀？可以给个示例case 和 编译命令吗？我按照cuda 的代码来编写仍然报错error: reference to global function 'child_launch' in global function，是需要添加什么编译参数吗？
--
楼主
0 点赞
0 帖子
楼主热帖
暂无数据常见问题新帖
暂无数据请问下，DTK-24.04.1版本的动态并行的功能如何使用呀？可以给个示例case 和 编译命令吗？我按照cuda 的代码来编写仍然报错error: reference to global function 'child_launch' in global function，是需要添加什么编译参数吗？
暂无数据 暂无数据
您好，若使用hipcc直接编译代码报“undefined reference to `__hipdrt_enable'”错误，请尝试分离编译链接操作：
hipcc -std=c++17 -fgpu-rdc -o dp.cpp.o -c dp.cpp
hipcc dp.cpp.o $ROCM_PATH/hip/lib/libgalaxyhipdrt.a -fgpu-rdc -o a.out
更完善的动态并行功能将于24.10提供。你好，DTK-24.04.1发版版本支持DP功能比较简单，也有较多受限情况，目前这项功能还处于继续开发过程中，预计DTK-24.10版本会对齐CUDA使用方式；
下边提供一个目前支持的简单case，仅供参考：
编译命令：
hipcc $ROCM_PATH/lib/libgalaxyhipdrt.a -fgpu-rdc test1.cpp
debug编译选项：-gdwarf-4
damon：
#include "hip/hip_runtime.h"
#include "hip/hip_device_runtime_api.h"
#define HIP_ENABLE_DRT
#define HIP_ENABLE_PRINTF
__device__ int error2_ = -1;
#define HIP_CHECK(stat) \
({ \
hipError_t restat = stat; \
if(restat != hipSuccess) \
{ \
printf("Error: hip error in line %d,stat:%d\n",__LINE__,restat) ; \
} \
})
#define REQUIRE(error) \
{ \
bool flag = error; \
if(!flag){ \
printf("error: in func: %s line: %d \n", __func__, __LINE__); \
} \
}
void CheckError2(int flag) {
int error = -1;
HIP_CHECK(hipStreamSynchronize(0));
HIP_CHECK(hipMemcpyFromSymbol(&error, error2_, sizeof(error)));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(error == flag);
}
void CheckError2() {
CheckError2(hipSuccess);
}
__device__ uint8_t u1_ = 0;
__device__ uint16_t u2_ = 0;
__device__ uint32_t u4_ = 0;
__device__ uint64_t u8_ = 0;
void CheckValue(uint8_t i_u1, uint16_t i_u2, uint32_t i_u4, uint64_t i_u8) {
uint8_t u1;
uint16_t u2;
uint32_t u4;
uint64_t u8;
HIP_CHECK(hipStreamSynchronize(0));
HIP_CHECK(hipMemcpyFromSymbol(&u1, u1_, sizeof(u1)));
HIP_CHECK(hipMemcpyFromSymbol(&u2, u2_, sizeof(u2)));
HIP_CHECK(hipMemcpyFromSymbol(&u4, u4_, sizeof(u4)));
HIP_CHECK(hipMemcpyFromSymbol(&u8, u8_, sizeof(u8)));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE((int) i_u1 == (int) u1);
REQUIRE((int) i_u2 == (int) u2);
REQUIRE((int) i_u4 == (int) u4);
REQUIRE((int) i_u8 == (int) u8);
}
__global__ void child_kernel_without_param() {
u1_ = -1;
u2_ = -1;
u4_ = -1;
u8_ = -1;
}
__global__ void child_kernel_with_param(uint8_t u1, uint16_t u2, uint32_t u4, uint64_t u8) {
u1_ = u1;
u2_ = u2;
u4_ = u4;
u8_ = u8;
}
__global__ void parent_kernel_multithread(bool has_param) {
if (has_param) {
hipLaunchKernelGGL_device("_Z23child_kernel_with_paramhtjm", dim3(10, 1), dim3(128, 1), 0, 0, (uint8_t) 1, (uint16_t) 2, (uint32_t) 3, (uint64_t) 4);
} else {
hipLaunchKernelGGL_device("_Z26child_kernel_without_paramv", dim3(10, 1), dim3(128, 1), 0, 0);
}
error2_ = __hipDpKernelTail();
}
__global__ void parent_kernel_null_stream(bool has_param) {
if (has_param) {
hipLaunchKernelGGL_device("_Z23child_kernel_with_paramhtjm", dim3(1, 1), dim3(1, 1), 0, 0, (uint8_t) 1, (uint16_t) 2, (uint32_t) 3, (uint64_t) 4);
} else {
hipLaunchKernelGGL_device("_Z26child_kernel_without_paramv", dim3(1, 1), dim3(1, 1), 0, 0);
}
error2_ = __hipDpKernelTail();
}
__global__ void parent_kernel(bool has_param) {
hipStream_t stream;
error2_ = hipStreamCreateWithFlags(&stream, hipStreamNonBlocking);
if (has_param) {
hipLaunchKernelGGL_device("_Z23child_kernel_with_paramhtjm", dim3(1, 1), dim3(1, 1), 0, stream, (uint8_t) 1, (uint16_t) 2, (uint32_t) 3, (uint64_t) 4);
} else {
hipLaunchKernelGGL_device("_Z26child_kernel_without_paramv", dim3(1, 1), dim3(1, 1), 0, stream);
}
error2_ = __hipDpKernelTail();
}
int main(){
{
printf("child_kernel_with_param\n");
{
parent_kernel<<<1, 1>>>(true);
CheckError2();
CheckValue(1, 2, 3, 4);
}
printf("child_kernel_without_param\n");
{
parent_kernel<<<1, 1>>>(false);
CheckError2();
CheckValue(-1, -1, -1, -1);
}
}
{
printf("child_kernel_with_param\n");
{
parent_kernel_null_stream<<<1, 1>>>(true);
CheckError2();
CheckValue(1, 2, 3, 4);
}
printf("child_kernel_without_param\n");
{
parent_kernel_null_stream<<<1, 1>>>(false);
CheckError2();
CheckValue(-1, -1, -1, -1);
}
}
printf("Unit_hipDynamicParallelism_deviceLaunch_parent_multithread\n");
{
parent_kernel<<<16, 256>>>(false);
parent_kernel<<<64, 256>>>(false);
CheckError2();
CheckValue(-1, -1, -1, -1);
}
printf("Unit_hipDynamicParallelism_deviceLaunch_child_multithread\n");
{
parent_kernel_multithread<<<128, 256>>>(false);
CheckError2();
CheckValue(-1, -1, -1, -1);
}
printf("Unit_hipDynamicParallelism_deviceLaunch_child_multithread\n");
{
parent_kernel_multithread<<<128, 256>>>(false);
CheckError2();
CheckValue(-1, -1, -1, -1);
}
return 0;
}