指针算术与共享内存

Pointer arithmetic with shared memory

本文关键字：共享 内存 指针      更新时间：2023-10-16

我不明白下面几行到底发生了什么:

  1. unsigned char *membershipChanged = (unsigned char *)sharedMemory;

  2. float *clusters = (float *)(sharedMemory + blockDim.x);

我假设在#1中sharedMemory有效地重命名为membershipChanged，但是为什么要将blockDim添加到sharedMemory指针。这个地址指向哪里?

sharedMemory是由extern __shared__ char sharedMemory[];创建的

我在CUDA kmeans实现中找到的代码。

void find_nearest_cluster(int numCoords,
                          int numObjs,
                          int numClusters,
                          float *objects,           //  [numCoords][numObjs]
                          float *deviceClusters,    //  [numCoords][numClusters]
                          int *membership,          //  [numObjs]
                          int *intermediates)
{
extern __shared__ char sharedMemory[];
//  The type chosen for membershipChanged must be large enough to support
//  reductions! There are blockDim.x elements, one for each thread in the
//  block.
unsigned char *membershipChanged = (unsigned char *)sharedMemory;
float *clusters = (float *)(sharedMemory + blockDim.x);
membershipChanged[threadIdx.x] = 0;
//  BEWARE: We can overrun our shared memory here if there are too many
//  clusters or too many coordinates!
for (int i = threadIdx.x; i < numClusters; i += blockDim.x) {
    for (int j = 0; j < numCoords; j++) {
        clusters[numClusters * j + i] = deviceClusters[numClusters * j + i];
    }
}
.....

sharedMemory + blockDim.x点距共享内存区域底部blockDim.x字节。

您可能这样做的原因是在共享内存中进行子分配。内核的启动站点(包括find_nearest_cluster)动态地为内核分配一定数量的共享存储。这段代码意味着两个逻辑上不同的数组驻留在sharedMemory——membershipChanged和clusters所指向的共享存储中。指针运算只是获取指向第二个数组的指针的一种方法。

相关文章：