Memcpy_async не работает с ролями конвейера ⇐ C++

[Anonymous]

Если я использую memcpy_async для каждого потока, все работает нормально, см. test_memcpy32 ниже.
Этот код предварительно извлекает данные в пределах одной деформации.
Я хочу расширить это, чтобы я мог предварительно выбирать данные в деформации 0 (блока) и использовать их (позже) в деформации 1 этого какого-либо блока. Для этого мне нужно использовать блочный конвейер.

Однако, если я использую memcpy_async с детализацией thread_block, я не могу заставить его работать, см. test_memcpy.

Я основывал свой код на документации конвейера CUDA.
Вот мой MCVE:

Код: Выделить всё

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include 
#include 
#include 
#include 

static constexpr auto consumer = cuda::pipeline_role::consumer;
static constexpr auto producer = cuda::pipeline_role::producer;

__global__ void test_memcpy32(int gdata[32]) {
assert(blockDim.x == 32);
__shared__ int sdata[32];
auto pipeline_warp = cuda::make_pipeline(); //per thread
pipeline_warp.producer_acquire(); //noop for per thread
const auto size = threadIdx.x & 1 ? sizeof(int): 0;
cuda::memcpy_async(&sdata[threadIdx.x], &gdata[threadIdx.x], size, pipeline_warp);
pipeline_warp.producer_commit();
//do other stuff
pipeline_warp.consumer_wait();
printf("tid: %i: gdata[%i] = %i, sdata[%i] = %i\n", threadIdx.x, threadIdx.x, gdata[threadIdx.x], threadIdx.x, sdata[threadIdx.x]);
pipeline_warp.consumer_release(); //noop for per thread

}

__global__ void test_memcpy(int gdata[32]) {
assert(blockDim.x == 64);
const auto group = cooperative_groups::this_thread_block();
__shared__ cuda::pipeline_shared_state PipelineState;
const auto warpid = threadIdx.x >= warpSize;
const auto laneid = threadIdx.x & 31;
//warp 0 = consumer, warp 1 = producer
const auto role = warpid ? producer : consumer;
auto pipeline_test = cuda::make_pipeline(group, &PipelineState, role);
__shared__ int sdata[32];
sdata[laneid] = 0;
const size_t size = threadIdx.x & 1 ? 4 : 4;
__syncthreads();
if (role == producer) {
pipeline_test.producer_acquire();
cuda::memcpy_async(group, &sdata[laneid], &gdata[laneid], size, pipeline_test);
pipeline_test.producer_commit();
} else {
assert(role == consumer);
pipeline_test.consumer_wait();
printf("tid: %i: gdata[%i] = %i, sdata[%i] = %i\n", threadIdx.x, laneid, gdata[laneid], laneid, sdata[laneid]);
pipeline_test.consumer_release();
}
}

int main() {
int* gdata;
int buffer[32];
cudaMalloc(&gdata, sizeof(buffer));
for (auto i = 0; auto& d: buffer) { d = i++; }
cudaMemcpy(gdata, buffer, sizeof(buffer), cudaMemcpyHostToDevice);
test_memcpy(gdata);  //does not work
test_memcpy32(gdata); //works just fine.
cudaDeviceSynchronize();
}

Код для каждого потока: test_memcpy32 выводит правильные данные, но код для каждого блока не работает. Он выводит все нули, хотя я следовал документации CUDA.
Я использую CUDA 13 и Visual Studio 17.9.6 на GTX 3070 (он же Compute 86).
Обновление версии до последней версии VS не имеет значения.
Обратите внимание, что для запуска этого кода требуется Ampere sm_80 или выше, поэтому он не работает на godbolt, потому что он использует sm_75 T4 GPU.
Если я запущу Compute-sanitizer с помощью инструмента синхронизации, я получу:

Код: Выделить всё

compute-sanitizer --tool=synccheck .\x64\Debug\MCVE_memcpy_async.exe
========= COMPUTE-SANITIZER
========= Barrier error detected.  Missing wait.
=========     at unsigned long long cuda::ptx::__4::mbarrier_arrive(unsigned long long *)+0x1b0 in mbarrier_arrive.h:21
=========     by thread (32,0,0) in block (0,0,0)
//Repeated for all threads in warp 1
=========     Barrier is located at shared address 0x0
=========         Device Frame: cuda::__4::barrier::arrive(long long)+0x9a0 in barrier_block_scope.h:130
=========         Device Frame: cuda::__4::pipeline::producer_commit()+0x6e0 in pipeline:270
=========         Device Frame: test_memcpy(int *)+0x1470 in kernel.cu:57
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: main in kernel.cu:72 [0x7a59] in MCVE_memcpy_async.exe

Что я делаю не так?

Подробнее здесь: https://stackoverflow.com/questions/797 ... line-roles