Câu hỏi 1: Tại sao việc khởi chạy kernel trong CUDA là bất đồng bộ?

• Khi một kernel được khởi chạy, CPU không đợi GPU hoàn thành.
• Bất đồng bộ cho phép CPU và GPU hoạt động song song.
• Để chờ kết quả, cần gọi hàm cudaDeviceSynchronize().
• Bất đồng bộ là cơ chế cốt lõi của CUDA để thực hiện nhiều tác vụ cùng lúc (như sao chép dữ liệu và tính toán).

kernel<<<grid, block>>>();
printf("CPU tiếp tục ngay lập tức!\n"); // CPU không đợi GPU
cudaDeviceSynchronize();  // Chờ đợi GPU hoàn thành

Câu hỏi 2: Sự khác biệt giữa cudaStreamSynchronize() và cudaDeviceSynchronize()?

• cudaDeviceSynchronize(): Đợi tất cả các tác vụ trong tất cả các luồng hoàn thành.
• cudaStreamSynchronize(stream): Chỉ đợi các tác vụ trong một luồng cụ thể hoàn thành.

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreate(&s2);

kernelA<<<grid, block, 0, s1>>>();
kernelB<<<grid, block, 0, s2>>>();

cudaStreamSynchronize(s1);  // Chỉ đợi A hoàn thành

Câu hỏi 3: Sự khác biệt giữa __syncthreads() và __threadfence()?

• __syncthreads(): Đồng bộ hóa thực thi các thread trong cùng một block.
• __threadfence(): Đảm bảo tính nhất quán của dữ liệu trong toàn bộ grid.

__global__ void kernel(int *data) {
    int idx = threadIdx.x;
    data[idx] = idx;
    __threadfence();   // Đảm bảo dữ liệu được ghi vào bộ nhớ toàn cục
    __syncthreads();   // Đồng bộ hóa các thread
}

Câu hỏi 4: Cách phân bổ bộ nhớ chia sẻ động trong kernel CUDA?

• Sử dụng từ khóa extern __shared__.
• Kích thước bộ nhớ chia sẻ được chỉ định bằng tham số thứ ba khi khởi chạy kernel.

__global__ void dynamicSharedMemory(float *input) {
    extern __shared__ float sharedData[];
    int idx = threadIdx.x;
    sharedData[idx] = input[idx] * 2;
    __syncthreads();
    input[idx] = sharedData[idx];
}

int main() {
    int N = 256;
    float *d_input;
    cudaMalloc(&d_input, N * sizeof(float));
    dynamicSharedMemory<<<1, N, N * sizeof(float)>>>(d_input);
}

Câu hỏi 5: CUDA hỗ trợ kernel mẫu như thế nào?

• Kernel có thể được định nghĩa dưới dạng mẫu C++.
• Compiler tạo ra phiên bản cụ thể của kernel khi gọi với toán tử <<>>.
• Dùng để hỗ trợ nhiều loại dữ liệu khác nhau.

template <typename T>
__global__ void templateKernel(T *a, T *b, T *c, int size) {
    int idx = threadIdx.x;
    c[idx] = a[idx] + b[idx];
}

int main() {
    templateKernel<float><<<1, 256>>>(d_a, d_b, d_c, 256);
    templateKernel<int><<<1, 256>>>(d_ai, d_bi, d_ci, 256);
}

Câu hỏi 6: Tại sao sử dụng cudaDeviceProp? Có thể lấy thông tin gì?

• Được sử dụng để truy vấn các thuộc tính của thiết bị.
• Một số thuộc tính quan trọng: name, multiProcessorCount, sharedMemPerBlock, maxThreadsPerBlock, warpSize, l2CacheSize.
• Giúp tối ưu hóa hiệu năng và lập lịch thuật toán.

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("GPU: %s, SMs=%d, Warp=%d\n", prop.name, prop.multiProcessorCount, prop.warpSize);

Câu hỏi 7: Khóa từ __restrict__ có chức năng gì?

• Thông báo cho trình biên dịch rằng con trỏ này không trùng lặp với các con trỏ khác.
• Cho phép trình biên dịch tối ưu hóa mạnh mẽ hơn.
• Thường được sử dụng đối với các tham số kernel quan trọng về hiệu năng.

__global__ void addition(const float *__restrict__ arr1,
                         const float *__restrict__ arr2,
                         float *__restrict__ result, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) result[idx] = arr1[idx] + arr2[idx];
}

Câu hỏi 8: CUDA thực hiện tối ưu hóa khởi chạy chậm như thế nào?

• CUDA lưu trữ yêu cầu khởi chạy kernel cho đến khi xảy ra các sự kiện như truyền dữ liệu, gọi cudaDeviceSynchronize(), hoặc luồng được đồng bộ hóa.
• Những kernel nhỏ có thể được tự động xử lý cùng lúc, cải thiện hiệu suất khởi chạy.

kernel1<<<1, 1>>>();
kernel2<<<1, 1>>>();
cudaDeviceSynchronize(); // Hai kernel được gửi cùng lúc tại đây

Câu hỏi 9: Tại sao GPU giỏi trong việc tính toán ma trận?

• Cấu trúc của GPU bao gồm nhiều đơn vị SIMD cùng băng thông cao.
• Tốt ở việc tính toán song song (áp dụng cùng một phép tính lên nhiều dữ liệu).
• Phép nhân ma trận phù hợp với mô hình này.
• Mỗi block thread xử lý một phần nhỏ của ma trận, tận dụng tính cục bộ dữ liệu.
• Bộ nhớ chia sẻ được sử dụng để lưu trữ các phần nhỏ, tăng tốc độ truy cập.

__global__ void matrixMultiply(const float *A, const float *B, float *C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0;
    for (int k = 0; k < N; ++k)
        sum += A[row * N + k] * B[k * N + col];
    C[row * N + col] = sum;
}

Câu hỏi 10: CUDA cấp phát bộ nhớ động trên thiết bị làm thế nào? Rủi ro gì?

• Từ CUDA 8 trở đi, hỗ trợ cấp phát bộ nhớ động trên thiết bị bằng malloc()/free().
• Có thể cấp phát bộ nhớ động bên trong kernel.
• Rủi ro: Chi phí cấp phát cao, dễ gây rạn vỡ bộ nhớ, thường không dùng trong các đường dẫn mã chạy thường xuyên.

__global__ void exampleAllocation() {
    int *ptr = (int*)malloc(sizeof(int));
    if (ptr) {
        *ptr = threadIdx.x;
        free(ptr);
    }
}

Câu hỏi 11: Lưu ý gì khi sử dụng __syncthreads() trong vòng lặp?

• Tất cả các thread phải thực hiện cùng số lần gọi __syncthreads().
• Nếu không, sẽ gây ra deadlock (các thread trong warp chờ đợi).
• Thường đặt __syncthreads() trong khối lệnh không điều kiện.

if (threadIdx.x < 16)
    __syncthreads(); // Sai: Chỉ một số thread thực hiện đồng bộ, gây deadlock

// Đúng:
if (threadIdx.x < 16)
    doSomething();
__syncthreads(); // Tất cả các thread đều thực hiện đồng bộ

Câu hỏi 12: Hàm thiết bị CUDA có thể gọi đệ quy không?

• Từ CUDA 5.0 (compute capability ≥ 3.5) trở đi, hỗ trợ gọi đệ quy.
• Tuy nhiên, chi phí hiệu năng cao (không gian ngăn xếp, thanh ghi, tràn).
• Mặc định, ngăn xếp nhỏ; có thể tăng kích thước ngăn xếp bằng cách sử dụng:

cudaDeviceSetLimit(cudaLimitStackSize, newSize);

__device__ int fibonacci(int n) {
    if (n <= 1) return n;
    return fibonacci(n-1) + fibonacci(n-2);
}

__global__ void exampleKernel() {
    printf("%d\n", fibonacci(5)); // In ra 5
}