标签: gpu-shared-memory

假设内核的一个线程尝试更新共享内存中的 4 个不同位置。如果任何其他线程覆盖了这些位置中的任何一个，我是否会导致该操作失败并被逆转？具体来说，这可以原子地执行吗？

mem[a] = x;
mem[b] = y;
mem[c] = z;
mem[d] = w;

我从matlab调用CUDA内核.

我之前被告知(David Kirk的书)每个线程只能占用16kb的共享内存,但我能够消耗更多:

__global__ void plain(float* arg)
{

    __shared__ float array[12000];
    int k;

    for (k=1;k<12000; k++)
    {
        array[k]=1;
    }   
}

CUDA C报告浮点数为4个字节,这意味着总数组大小为48Kb,大于12Kb.它运行良好,那怎么可能呢？

我也被告知在 GPU中共享内存的大小非常小 - 我该怎么办呢？ 每个块的最大共享内存很重要.我的卡的每块最大共享内存是49152字节,但我能够运行上面的代码,每块1000个线程.

似乎每块会使用49Kb,这是不对的.是SM一次只为一个块提供服务而在dong中保留了每个线程块只能有49Kb的条件吗？

每个块的49Kb共享内存如何与每个线程的16Kb共享内存协调？

谢谢

我正在尝试在 CUDA 中实现并行合并算法。该算法被设计为在一个线程块中执行。其基本思想是计算两个输入序列中每个元素的全局排名。由于这两个输入序列已排序，因此元素的全局排名等于其在原始序列中的索引加上其在通过二分搜索计算的另一个序列中的排名。我认为实现这种算法的最佳策略是将两个序列加载到共享内存中以减少全局内存读取。然而，当我比较两个版本的实现时，一种使用共享内存，一种不使用共享内存，我看不到性能的提高。我想知道我是否做错了什么。

硬件：GeForce GTX 285、Linux x86_64。

两种实现方式合并两个 1024 个元素的序列的时间约为 0.068672 ms。

__global__ void localMerge(int * A, int numA,int * B,int numB,int * C){
    extern __shared__ int  temp[]; // shared memory for A and B; 
    int tx=threadIdx.x;
    int size=blockDim.x;
    int *tempA=temp;
    int *tempB=temp+numA;
    
    int i,j,k,mid;
    //read sequences into shared memory 
    for(i=tx;i<numA;i+=size){
        tempA[i]=A[i];
    }
    for(i=tx;i<numB;i+=size){
        tempB[i]=B[i];
    }
    __syncthreads();
    //compute global rank for elements in sequence A
    for(i=tx;i<numA;i+=size){
        j=0;
        k=numB-1;
        if(tempA[i]<=tempB[0]){
            C[i]=tempA[i];
        }
        else if(tempA[i]>tempB[numB-1]){
            C[i+numB]=tempA[i];
        }
        else{
            while(j<k-1){
                mid=(j+k)/2;
                if(tempB[mid]<tempA[i]){
                    j=mid; …

我有以下代码，它使用共享内存执行平铺矩阵转置以提高性能。共享内存用 1 列填充，以避免 32x32 线程块的存储体冲突。

__global__ void transpose_tiled_padded(float *A, float *B, int n)
{
    int i_in = blockDim.x*blockIdx.x + threadIdx.x;
    int j_in = blockDim.y*blockIdx.y + threadIdx.y;
    int i_out = blockDim.x*blockIdx.y + threadIdx.x;
    int j_out = blockDim.y*blockIdx.x + threadIdx.y;

    extern __shared__ float tile[];

    // coalesced read of A rows to (padded) shared tile column (transpose)
    tile[threadIdx.y + threadIdx.x*(blockDim.y+1)] = A[i_in + j_in*n];
    __syncthreads();

    // coalesced write from (padded) shared tile column to B rows
    B[i_out + j_out*n] = tile[threadIdx.x + threadIdx.y*(blockDim.x+1)]; …

我对 CUDA 编程相当陌生，我正在尝试编写一个 CUDA 内核，用于仅对 3 维张量的 1 维进行并行缩减，该张量是float馈入内核的行主展平数组。

换句话说，我试图用,和numpy.sum的有限轴组合重写。axis=0axis=1axis=2

我已经成功实现了“reduce over axis 0”和“reduce over axis 1”，但是“reduce over axis2”的性能问题让我在这里发布了一个问题来寻求建议。

内核以一维网格和一维块配置启动，并将每个线程映射到缩减输出张量的每个元素。所以，它应该是这样的：

这是我的内核：

__global__ void kernel_reduce_sum_3d_try02(
        float* g_idata,
        float* g_odata,
        int dim0,
        int dim1,
        int dim2,
        int overaxis0,
        int overaxis1,
        int overaxis2)
{

    if (overaxis0 == 0 && overaxis1 == 0 && overaxis2 == 1) { 
        // static shared memory
        __shared__ float smem_store[BLOCK_SIZE];

        // set thread ID
        //unsigned int tid = threadIdx.x;
        unsigned int tid …

我正在编写一个 CUDA 程序，它在共享内存中定义了一个数组。我需要做的是只允许一个线程写入此数组中的每个索引，即到达此写入指令的第一个线程应更改其值，但同一扭曲或下一个扭曲中的任何其他线程应读取写入的值。

这是代码片段：

char* seq_copied = seqs + (njobNew * halfLength); //this is the shared memory array
if (seq_copied[seq_1_index] == false) { //here is the condition that I need to check with only one thread
    seq_copied[seq_1_index] = true; //and this is the write that should be written by only one thread
    printf("copy seq_shared seq_1_index = %d,  block = %d \n", seq_1_index, blockIdx.x);
}

现在发生的情况是，warp 中的所有线程都执行这些确切的指令序列，因此 if 条件中的剩余代码被执行了 32 次。我只需要执行一次。

我怎样才能做到这一点？

在 CUDA 编程指南的共享内存部分中，它指出 warp 的共享内存访问不是序列化的，而是广播读取的。

然而，它没有说明如果整个块请求相同的内存地址会发生什么。warp 之间的访问是串行的还是 CUDA 可以广播到整个块。

我的案例的演示代码

// Assume 1024 sized int array
__global__ add_from_shared(int* i, int* j, int* out)
{
    __shared__ int shmem[1024];
    shmem[threadIdx.x] = i[threadIdx.x];
    ...
    Do some stuff
    ...
    // Is the shared memory call here serilized between warps or is it a broadcast over the entire block?
    j[threadIdx.x] += shmem[0];
}  

谢谢

我有时会在 CUDA 内核中看到以下共享内存声明，但我不确定它的含义：

extern __shared__ T shmem[][SZ]

作为SZ编译时常量。内核启动如下：

kernel<<<grid, block, shared_memory_size>>>()

我的问题是：

• 为什么动态和静态共享内存的混合会很有用（除了简化程序员的地址计算之外）？
• 共享内存缓冲区 shmem 的总大小是多少？是吗shared_memory_size * SZ？
• 与问题2相关：假设我可以在编译时计算shared_memory_size，我将如何重写共享内存声明以使其成为静态？shared T shmem[SZ*shared_memory_size]？

我正在编写一些 N 体模拟代码，在 CUDA 中针对 Volta 和图灵系列卡进行短程交互。我计划使用共享内存，但我不太清楚这样做时如何避免银行冲突。由于我的交互是本地的，我计划将我的粒子数据分类到本地组中，我可以将这些数据发送到每个 SM 的共享内存（还没有担心粒子的邻居正在从另一个 SM 工作。为了变得更好性能（避免库冲突），仅每个线程从/向共享内存的不同地址读取/写入就足够了，但每个线程可以无序访问该内存而不会受到惩罚？

我看到的所有信息似乎只提到内存被合并以从全局内存到共享内存的复制，但我没有看到任何关于扭曲（或整个 SM）中的线程是否关心共享内存中的合并。