最简示例 简介洗牌函数 之 __shfl_sync() cuda 之 shuffle

虽然叫做洗牌函数，但是效果只是跟真正的洗牌有点像，但还是不一样。扑克牌不可以复制，但这里的shuffle函数很多时候会复制牌值。

这个函数具有广播功能，当同一个warp内的其他线程想获得某一个特定线程中的某个局部变量的值时，使用__shfl_sync()。

__shfl_sync(0xffffffff, value, 9, 32);//把 laneId == 9线程的 value变量的值， return 广播给其他线程

#include <cuda_runtime.h>
//#include <iostream>

#include <stdio.h>
#define warpSize 32

__global__ void bcast(float* a, float* b) {
    int laneId = threadIdx.x & 0x1f;
    float value;

    value = a[laneId];
    value = __shfl_sync(0xffffffff, value, 9, 32);//把 laneId == 9线程的value变量的值， return 广播给其他线程
    b[laneId] = value;

}

void printVector(char* desc, float* ptr_vec, unsigned int n){
    printf("%s =\n", desc);

    for(int i=0; i<n; i++){
        printf(" %5.2f ",ptr_vec[i]);
    }

    printf("\n");
}

int main() {

    float* a_h = NULL;
    float* a_d = NULL;
    float* b_h = NULL;
    float* b_d = NULL;

    a_h = (float*)malloc(warpSize*sizeof(float));
    b_h = (float*)malloc(warpSize*sizeof(float));

    for(int i=0; i<warpSize; i++){
        a_h[i] = i+100.0;
    }

    for(int i=0; i<warpSize; i++){
        b_h[i] = i+100;
    }

    printVector("a_h",a_h, warpSize);
    printVector("b_h",b_h, warpSize);

    cudaMalloc((void**)&a_d, warpSize*sizeof(float));
    cudaMalloc((void**)&b_d, warpSize*sizeof(float));

    cudaMemcpy(a_d, a_h, warpSize*sizeof(float), cudaMemcpyHostToDevice);  
    cudaMemcpy(b_d, b_h, warpSize*sizeof(float), cudaMemcpyHostToDevice);  

    bcast<<< 1, warpSize >>>(a_d, b_d);
    cudaDeviceSynchronize();

    cudaMemcpy(b_h, b_d, warpSize*sizeof(float), cudaMemcpyDeviceToHost);

    printVector("b_d", b_h, warpSize);

    cudaFree(a_d);
    cudaFree(b_d);

    return 0;
}