我将要问的问题也有类似的问题，但是我觉得这些问题都不是我真正寻找的问题的核心。我现在所拥有的是一种CUDA方法，该方法要求在共享内存中定义两个数组。现在，数组的大小由执行开始后读取到程序中的变量给定。因此，由于定义共享数组的大小需要在编译时知道该值，因此我无法使用该变量来定义数组的大小。我不想做这样的事情，__shared__ double arr1[1000]因为手动输入大小对我来说是没有用的，因为这会根据输入而改变。同样，我不能用来#define为尺寸创建一个常数。
现在，我可以按照类似于手册（http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared）中的示例进行操作，例如
extern __shared__ float array[];
__device__ void func() // __device__ or __global__ function
{
short* array0 = (short*)array;
float* array1 = (float*)&array0[128];
int* array2 = (int*)&array1[64];
}
但这仍然是一个问题。根据我的阅读，定义共享数组总是使内存地址成为第一个元素。这意味着我需要使第二个数组移位第一个数组的大小，如本例中所示。但是第一个数组的大小取决于用户输入。
另一个问题（Cuda Shared Memory数组变量）也有类似的问题，他们被告知创建一个单独的数组作为两个数组的数组，并简单地调整索引以正确匹配数组。尽管这确实可以满足我的要求，但看起来却很混乱。有什么办法可以解决，以便我仍然可以维护两个独立的数组，每个数组的大小都由用户定义为输入？
当将动态共享内存与CUDA一起使用时，只有一个指针传递给内核，该指针以字节为单位定义请求/分配区域的开始：
extern __shared__ char array[];
没有办法以不同的方式处理它。但是，这不会阻止您拥有两个用户大小的阵列。这是一个可行的示例：
$ cat t501.cu
#include <stdio.h>
__global__ void my_kernel(unsigned arr1_sz, unsigned arr2_sz){
extern __shared__ char array[];
double *my_ddata = (double *)array;
char *my_cdata = arr1_sz*sizeof(double) + array;
for (int i = 0; i < arr1_sz; i++) my_ddata[i] = (double) i*1.1f;
for (int i = 0; i < arr2_sz; i++) my_cdata[i] = (char) i;
printf("at offset %d, arr1: %lf, arr2: %d\n", 10, my_ddata[10], (int)my_cdata[10]);
}
int main(){
unsigned double_array_size = 256;
unsigned char_array_size = 128;
unsigned shared_mem_size = (double_array_size*sizeof(double)) + (char_array_size*sizeof(char));
my_kernel<<<1,1, shared_mem_size>>>(256, 128);
cudaDeviceSynchronize();
return 0;
}
$ nvcc -arch=sm_20 -o t501 t501.cu
$ cuda-memcheck ./t501
========= CUDA-MEMCHECK
at offset 10, arr1: 11.000000, arr2: 10
========= ERROR SUMMARY: 0 errors
$
如果您随机排列混合数据类型的数组，则需要手动对齐数组起点（并请求足够的共享内存），或者使用对齐指令（并确保请求足够的共享内存），或者使用有助于对齐的结构。
本文收集自互联网，转载请注明来源。
如有侵权，请联系 [email protected] 删除。
我来说两句