What will happen if I mistakely use __align__ on shared memory?

[202476410arsmart]

Hi! I am learning other people’s code and have this within it:

	__shared__ __align__(4 * 1024) char smem[6 * 1024]; 

	float *ashare = reinterpret_cast<float *>(smem);
	float *bshare =
		reinterpret_cast<float *>(smem + 4 * 1024);
	float *cshare =
		reinterpret_cast<float *>(smem + 3 * 1024);

align to 41024!!! But later I need to use a cshare locate at smem+31024 and seems to break the previous align…

I am not sure what will happen inside? The compiler will change my cshare back to smem+41024? Or still use it as +3smem but very! slow?

Thank you!!!

[Robert_Crovella]

@202476410arsmart

__shared__ __align__(4 * 1024) char smem[6 * 1024];

That says: "align the pointer you give to me so that it starts on a 4*1024 byte boundary. No problem there. For example, pretend CUDA set smem to 0. It would fit that request.

@202476410arsmart

float *ashare = reinterpret_cast<float *>(smem);

That says, “take the pointer that was aligned on a 4*1024 byte boundary, and alias another float pointer to point to it.” No problem that I can see. You can access an array of floats starting on such a boundary. Using smem of 0, that means this pointer would start at address 0.

@202476410arsmart

	float *bshare =
		reinterpret_cast<float *>(smem + 4 * 1024);

That says “take the pointer that was aligned on a 4*1024 byte boundary, add 4*1024 to it, and alias another float pointer to point to it.” No problem that I can see. Using smem of 0, that means this pointer would point to address 4096. It would be perfectly acceptable to have a float array whose starting address was 4096.

@202476410arsmart

float *cshare =
		reinterpret_cast<float *>(smem + 3 * 1024);

That says “take the pointer that was aligned on a 4*1024 byte boundary, add 4*1024 to it, and alias another float pointer to point to it.” No problem that I can see. Using smem of 0, this pointer would point to 3072. It would be perfectly acceptable to have a float array whose starting address was 3072.

[202476410arsmart]

Thank you!! I realized, align is for memory’s allocation in the disk, but my array setting is to use the memory, of which the two things are irrelevant! I fully understand your answer now!!!