Câu 1: Kiến trúc SIMT trong GPU hoạt động như thế nào?

Giải thích:

• CPU sử dụng mô hình SIMD, GPU dùng SIMT với các luồng độc lập
• Warp (32 luồng) là đơn vị thực thi cơ bản
• Phân nhánh trong warp gây warp divergence làm giảm hiệu suất

Ví dụ minh họa:

__global__ void simtExample(int *data) {
    int idx = threadIdx.x;
    if (idx % 3 == 0) {
        data[idx] += 5;
    } else {
        data[idx] -= 2;
    }
}

Câu 2: Xung đột ngân hàng bộ nhớ chia sẻ (Bank Conflict) và cách xử lý?

Vấn đề:

• 32 ngân hàng bộ nhớ chia sẻ có thể gây xung đột khi truy cập đồng thời
• Xung đột buộc truy cập tuần tự làm giảm hiệu suất

Giải pháp:

1. Thêm độ dãn (padding) cho mảng
2. Thiết kế mẫu truy cập phân tán
3. Sử dụng __shared__ float data[32][34] thay vì 32x32

__global__ void bankConflict(float *out) {
    __shared__ float buffer[32][34];
    int tid = threadIdx.x;
    buffer[tid][tid] = tid * 1.5f;
    out[tid] = buffer[tid][tid];
}

Câu 3: Tràn thanh ghi (Register Spilling) và ảnh hưởng?

Nguyên nhân:

• Giới hạn ~255 thanh ghi/luồng
• Tràn sang bộ nhớ địa phương (Local Memory) trên VRAM

Phương pháp kiểm tra:

nvcc --ptxas-options=-v -arch=sm_75 main.cu

__global__ void spillExample() {
    float temp[512];
    temp[threadIdx.x] = threadIdx.x * 0.1f;
}

Câu 4: Chỉ số Occupancy và vai trò trong tối ưu hiệu suất

Công thức: Occupancy = (Số warp hoạt động thực tế)/(Số warp tối đa)

Yếu tố ảnh hưởng:

• Số lượng thanh ghi/lượt
• Kích thước block
• Dung lượng bộ nhớ chia sẻ

int blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, kernel, 0, 0);

Câu 5: Phân biệt các chỉ thị CUDA

Chỉ thị	Nơi định nghĩa	Nơi thực thi	Đối tượng gọi
__host__	Hàm CPU	CPU	CPU
__device__	Hàm GPU	GPU	GPU
__global__	Hàm kernel	GPU	CPU

__device__ int multiply(int a, int b) { return a * b; }
__global__ void calc(int *result) { result[threadIdx.x] = multiply(3, 4); }

Câu 6: Bộ nhớ thống nhất (Unified Memory) và hạn chế

Tính năng:

• CUDA 6.0 trở lên hỗ trợ cudaMallocManaged()
• Tự động di chuyển dữ liệu giữa CPU/GPU

Thách thức:

• Trễ do di chuyển trang (page migration)
• Tối ưu bằng cudaMemPrefetchAsync()

int *data;
cudaMallocManaged(&data, 1024*sizeof(int));
parallelKernel<<<16,64>>>(data);
cudaDeviceSynchronize();

Câu 7: Đặc điểm kiến trúc GPU

Ưu điểm: Hàng ngàn luồng song song che giấu độ trễ

Hạn chế: Hiệu suất luồng đơn thấp do phần cứng đơn giản

__global__ void latencyHiding(float *array) {
    int i = threadIdx.x;
    for(int j=0; j<1000; j++) {
        array[i] = sqrt(array[i]) + 1.0f;
    }
}

Câu 8: Lệnh Warp Shuffle trong tính toán song song

Ứng dụng: Thuật toán giảm tổng (reduction), quét (scan)

__global__ void warpSum(float *input) {
    float value = input[threadIdx.x];
    for(int offset=16; offset>0; offset/=2) {
        value += __shfl_down_sync(0xFFFFFFFF, value, offset);
    }
    if(threadIdx.x % 32 == 0) input[threadIdx.x] = value;
}

Câu 9: Quản lý độ ưu tiên luồng (Stream Priority)

int lowPriority, highPriority;
cudaDeviceGetStreamPriorityRange(&lowPriority, &highPriority);

cudaStream_t streamA, streamB;
cudaStreamCreateWithPriority(&streamA, cudaStreamDefault, highPriority);
cudaStreamCreateWithPriority(&streamB, cudaStreamDefault, lowPriority);

Câu 10: Vai trò của căn chỉnh bộ nhớ

Phương pháp: Sử dụng cudaMallocPitch() cho mảng 2D

float *deviceData;
size_t pitch;
cudaMallocPitch(&deviceData, &pitch, width*sizeof(float), height);

Câu 11: So sánh bộ nhớ chia sẻ và thanh ghi

Tiêu chí	Thanh ghi	Bộ nhớ chia sẻ
Tốc độ	~1TB/s	~100GB/s
Phạm vi	Luồng	Khối
Dung lượng	~256/luồng	48KB/khối

Câu 12: Ảnh hưởng băng thông PCIe

Giải pháp:

• Sử dụng bộ nhớ khóa trang (Pinned Memory)
• Truyền dữ liệu bất đồng bộ
• Tối ưu hóa luồng tính toán

float *hostData, *deviceData;
cudaMallocHost(&hostData, N*sizeof(float));
cudaMalloc(&deviceData, N*sizeof(float));
cudaMemcpyAsync(deviceData, hostData, N*sizeof(float), cudaMemcpyHostToDevice);

Câu 13: Xử lý lỗi CUDA

#define CHECK_CUDA(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA Error: %s at %s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__); \
        } \
    } while(0)

Câu 14: Thuật toán quét song song (Prefix Sum)

__global__ void prefixSum(int *data) {
    __shared__ int cache[256];
    int tid = threadIdx.x;
    cache[tid] = data[tid];
    __syncthreads();
    
    for(int step=1; step<256; step*=2) {
        int val = 0;
        if(tid >= step) val = cache[tid - step];
        __syncthreads();
        cache[tid] += val;
        __syncthreads();
    }
    data[tid] = cache[tid];
}