感谢 @ANYA 在我[上一篇提问](https://forum.sourcefind.cn/thread/50a90b01-8e22-11f0-bd66-0242ac150005)中告诉我 >老师您好，全部用hip写的代码，`ds_read` `mmac`的等待不需要人为去加，编译器会去加上的，只需要在`ds_write`后加个同步，保证数据写道lds中；如果代码中使用内嵌汇编来处理lds/global相关的操作，那么等待指令需要自己人为的加上，编译器不会处理,global的读取需要`vmcnt`来等待，lds的读取需要`lgkmcnt`来等待 我在那个问题中使用的`PREFETCH_LDS`是一个封装了`ds_read`内联汇编的宏，故而需要手动添加特殊同步指令 ```c++ #define PREFETCH_LDS(a, offset_a) \ asm volatile("ds_read_m32x8_b32 %0, %1 offset:0\n\t" : "+v"(a.f4), "+v"(offset_a)); ``` 现在我想了解一下老师所说的**用hip写的`dsread`代码** 经过查找资料，我认为这里说的**用hip写的`dsread`代码**指的是`__builtin_amdgcn_ds_read_*`编译器内置函数(**如果不是请告诉我**)。以`__builtin_amdgcn_ds_read_m32x8f32`为例，这个函数的第一个参数表示LDS地址，类型需要为`__local float *`，第二个参数表示地址偏移，类型需要为`const short`. 问题出现在第二个参数上：每个线程mma时需要的数据不同，那么这个偏移必定是一个依赖于`threadIdx.x`的表达式，从而不能是`constexpr`或`const` 我下面的代码反映了这个问题： ```c++ for (int r = 0; r < 2; r++) { int input_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; const int input_lds_addr = input_mma_load_inner_offset + input_mma_load_offset; int input_lds_addr_in_byte = input_lds_addr * sizeof(float); // PREFETCH_LDS(_b, input_lds_addr_in_byte); _b = __builtin_amdgcn_ds_read_m32x8f32(&smem[0], input_lds_addr); int weight_mma_load_offset = ((tid % 64) / 8 + r * 8) * 64 + (tid % 8) * 4; const int weight_lds_addr = weight_mma_load_inner_offset + weight_mma_load_offset + weight_dsread_offset; int weight_lds_addr_in_byte = weight_lds_addr * sizeof(float); // PREFETCH_LDS(_a, weight_lds_addr_in_byte); _a = __builtin_amdgcn_ds_read_m32x8f32(&smem[0], weight_lds_addr); // ... // mmac computation after that... } ``` 这时编译报错 ```c++ /root/private_data/conv3d/src/conv3d/TensorConv3d_64x64x16_alt.cpp:158:18: error: argument to '__builtin_amdgcn_ds_read_m32x8f32' must be a constant integer _b = __builtin_amdgcn_ds_read_m32x8f32(&smem[0], input_lds_addr); ^ ~~~~~~~~~~~~~~ /root/private_data/conv3d/src/conv3d/TensorConv3d_64x64x16_alt.cpp:164:18: error: argument to '__builtin_amdgcn_ds_read_m32x8f32' must be a constant integer _a = __builtin_amdgcn_ds_read_m32x8f32(&smem[0], weight_lds_addr); ^ ~~~~~~~~~~~~~~~ 2 errors generated when compiling for gfx928. ``` 请问`__builtin_amdgcn_ds_read_*`正确的使用方式是怎样的？有相应的文档和参考代码吗？