2012-09-22 125 views
2

是否有任何应用程序级别的API可用于由CUDA中的CTA分配的空闲共享内存?我想重新使用我的CTA来完成另一项任务,在开始执行该任务之前,我应该清除先前任务使用的内存。CUDA中的免费共享内存

+0

共享内存不会在内核调用之间持续存在:您的问题是什么?总的来说,启动一个新内核要比在同一内核中完成更多(异构)任务要好得多。 –

+0

我想为其他任务使用相同的CTA,因此我打算与__syncThreads()同步并将其用于其他任务。 –

+0

Stefano,内核融合是在CUDA中分摊内核启动开销的早期方式。 – ArchaeaSoftware

回答

4

共享内存在内核启动时静态分配。您可以选择在内核中指定一个未施胶共享分配:

__global__ void MyKernel() 
{ 
    __shared__ int fixedShared; 
    extern __shared__ int extraShared[]; 
    ... 
} 

第三内核启动参数,然后指定多少共享内存如何对应于未分级的分配。

MyKernel<<<blocks, threads, numInts*sizeof(int)>>>(...); 

分配给内核启动共享存储器的总量是在内核中声明的量,再加上共享存储器内核参数,再加上对准开销的总和。你不能“释放”它 - 它在内核启动期间保持分配。

对于通过执行的多个阶段去需要使用共享内存用于不同目的的内核,你可以做的是重用共享内存指针的内存 - 在未分级的声明中使用指针运算。

喜欢的东西:

__global__ void MyKernel() 
{ 
    __shared__ int fixedShared; 
    extern __shared__ int extraShared[]; 
    ... 
    __syncthreads(); 
    char *nowINeedChars = (char *) extraShared; 
    ... 
} 

我不知道用这个成语任何SDK的样本,虽然threadFenceReduction示例声明一个__shared__ bool并且还使用共享内存来存放减少的部分和。

+1

内核融合可以很好地工作,但要注意注册使用情况。为内核分配的寄存器数量在内核启动时也是静态的,因此,较大内核中有较高寄存器要求的任何部分都会相应地减少整个内核的占用率。 – ArchaeaSoftware