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

[Anonymous]

Если я использую memcpy_async для каждого потока, все работает нормально, см. test_memcpy32 ниже.
Однако, если я использую 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 не имеет никакого значения.
Что я делаю не так?

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