4
数学库
dsread编译报错
2025-09-23 16:10:07
95
在使用`__builtin_amdgcn_ds_read_m32x8f32`时报错。代码如下 ```c++ typedef float native_float4 __attribute__((__vector_size__(16))); typedef float native_float2 __attribute__((__vector_size__(8))); union SmemData { native_float4 f4; native_float2 f2[2]; }; SmemData _a[2], _b[2]; // .... // load lds data to registers for further mma { int next_mma_offset = load_flag * 64 * 16; int _weight_addr = ((tid % 64) / 8 + 8) * 64 + (tid % 8) * 4; int next_mma_weight_addr = next_mma_offset + _weight_addr + weight_smem_offset + weight_mma_load_cross_offset; int next_mma_weight_addr_in_byte = next_mma_weight_addr * sizeof(float); // PREFETCH_LDS(_a[1], next_mma_weight_addr_in_byte); _a[1].f4 = __builtin_amdgcn_ds_read_m32x8f32(&smem[next_mma_weight_addr], 0); int _input_addr = ((tid % 64) / 8 + 8) * 64 + (tid % 8) * 4; int next_mma_input_addr = next_mma_offset + _input_addr + input_smem_offset + input_mma_load_cross_offset; int next_mma_input_addr_in_byte = next_mma_input_addr * sizeof(float); // PREFETCH_LDS(_b[1], next_mma_input_addr_in_byte); _b[1].f4 = __builtin_amdgcn_ds_read_m32x8f32(&smem[next_mma_input_addr], 0); } ``` 编译后报错为 ```c++ [ 38%] Building CXX object CMakeFiles/conv3d_DoubleBufferTensor64x64x16_ds_read.dir/src/conv3d/DoubleBufferTensorConv3d_64x64x16_ds_read.cpp.o Call parameter type does not match function signature! i32 0 i16 %208 = call <4 x i32> @llvm.amdgcn.ds.read.m32x8b32.v4i32.p3f32(float addrspace(3)* %207, i32 0), !dbg !2058 Call parameter type does not match function signature! i32 0 i16 %211 = call <4 x i32> @llvm.amdgcn.ds.read.m32x8b32.v4i32.p3f32(float addrspace(3)* %210, i32 0), !dbg !2061 fatal error: error in backend: Broken function found, compilation aborted! dcc: error: clang frontend command failed with exit code 70 (use -v to see invocation) clang version 15.0.0 Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /opt/dtk-24.04.2/llvm/bin dcc: note: diagnostic msg: Error generating preprocessed source(s). make[2]: *** [CMakeFiles/conv3d_DoubleBufferTensor64x64x16_ds_read.dir/build.make:90: CMakeFiles/conv3d_DoubleBufferTensor64x64x16_ds_read.dir/src/conv3d/DoubleBufferTensorConv3d_64x64x16_ds_read.cpp.o] Error 1 make[1]: *** [CMakeFiles/Makefile2:272: CMakeFiles/conv3d_DoubleBufferTensor64x64x16_ds_read.dir/all] Error 2 make: *** [Makefile:91: all] Error 2 ``` 我明明用的是`float4`,为什么还是`i16`?请问如何解决,该如何正确调用?
dcu-k100-ai显卡运行一段时间后
全部评论(2)
只看楼主
C
changhl
2楼
老师,您好: 1.内部使用dtk25041复现,建议您后续可以基于最新的dtk开发及编译。 2.对于上面的dsread编译bug,为内置函数调用错误,__builtin_amdgcn_ds_read_m32x8f32 传参需要转为lds指针,参考案例代码第25行
2025/9/23
网
网友DvlKdg
1楼
1. 报错Call parameter type does not match function signature! __builtin_amdgcn_ds_read_m32x8f32第一个参数要强制转换为lds指针类型
2025/9/23
没有更多内容了
--
楼主
Level.2
0 帖子
楼主热帖
暂无数据