Tại sao các biến thành viên không được chia sẻ?

Question

Tôi muốn khởi tạo một lớp trong mã CUDA, chia sẻ một số thành viên của nó với các chủ đề khác trong cùng một khối.

Tuy nhiên, khi cố gắng biên dịch mã sau, tôi nhận được lỗi: »thuộc tính" được chia sẻ "không áp dụng ở đây« (nvcc phiên bản 4.2).

class SharedSomething { 

public: 
    __shared__ int i; // this is not allowed 
}; 

__global__ void run() { 

    SharedSomething something; 
} 

Lý do đằng sau điều đó là gì? Có một công việc xung quanh để đạt được hành vi mong muốn (các thành viên được chia sẻ của một lớp trên một khối) không?

Answer 1

Rost giải thích lý do đằng sau giới hạn. Để trả lời phần thứ hai của câu hỏi, một cách giải quyết đơn giản là để hạt nhân khai báo bộ nhớ dùng chung và khởi tạo một con trỏ đến nó do lớp đó sở hữu, ví dụ: trong hàm tạo lớp. Thí dụ.

class Foo 
{ 
public: 
    __device__ 
    Foo(int *sPtr) : sharedPointer(sPtr, gPtr) { 
    sharedPointer[threadIdx.x] = gPtr[blockIdx.x * blockDim.x + threadIdx.x]; 
    __syncthreads(); 
    } 

    __device__ 
    void useSharedData() { printf("my data: %f\n", sharedPointer[threadIdx.x]); } 

private: 
    int *sharedPointer; 
}; 

__global__ void example(int *gData) 
{ 
    __shared__ int sData[BLOCKDIM]; 

    Foo f(sData, gData); 

    f.useSharedData(); 
} 

Caveat: Mã văn bản trong trình duyệt, chưa được xác minh, chưa được kiểm tra (và đó là một ví dụ nhỏ, nhưng khái niệm này kéo dài đến mã thực — Tôi đã sử dụng kỹ thuật này bản thân mình).

Answer 2

Đối tượng được đánh dấu là __shared__ nằm trong bộ nhớ dùng chung được dành riêng cho mỗi khối chuỗi. Nó có kích thước giới hạn và có cùng tuổi thọ như khối chuỗi.

Vì vậy, đây là lý do tại sao bạn không thể khai báo thành viên của lớp là được chia sẻ - cuộc đời của họ không được quản lý bởi cá thể lớp, mà theo khối chuỗi. Có thể static thành viên lớp học có thể được chia sẻ nhưng không kiểm tra.

Xem CUDA Programming Guide để biết chi tiết, phần B.2.3.