Fixed the issues of including Thrust and MIOpen

Author: coketaste

docker/rocm/Dockerfile

@@ -1,7 +1,7 @@
 # GPU Programming 101 - ROCm Development Container
 # Based on AMD's official ROCm development image - used as-is for maximum compatibility
 
-FROM rocm/dev-ubuntu-24.04:latest
+FROM rocm/dev-ubuntu-24.04:7.0-complete
 
 # Metadata
 LABEL maintainer="GPU Programming 101"

modules/module1/examples/03_matrix_multiplication_hip.cpp

@@ -80,8 +80,8 @@ __global__ void matrixMultiplyAMDOptimized(float *A, float *B, float *C, int N)
 
     float sum = 0.0f;
 
-    // Unrolled tile loop for better instruction scheduling
-    #pragma unroll 4
+    // ROCm 7: Use clang loop optimization hints instead of fixed unroll count
+    // The compiler will determine optimal unrolling based on target architecture
     for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; t++) {
         // Coalesced loads
         tileA[hipThreadIdx_y][hipThreadIdx_x] = 

modules/module1/examples/Makefile

@@ -33,8 +33,8 @@ HIP_DEBUG_FLAGS = -std=c++17 -g
 
 # ROCm 7 note: hipcc may require explicit --rocm-path when HIP runtime/version files
 # aren't in legacy locations. See ROCm docs on file structure reorg.
-# Prefer environment ROCM_PATH if set; fall back to /opt/rocm.
-ROCM_PATH ?= /opt/rocm
+# Prefer environment ROCM_PATH if set; prefer ROCm 7.0.0, fall back to any ROCm installation.
+ROCM_PATH ?= $(shell ls -d /opt/rocm-7.0.0 2>/dev/null || ls -d /opt/rocm* 2>/dev/null | head -1 || echo /opt/rocm)
 # Prefer hipconfig (ROCm 7) to locate ROCm root, then fall back to hipcc path.
 HIPCONFIG_BIN := $(shell command -v hipconfig 2>/dev/null)
 ifneq ($(HIPCONFIG_BIN),)

modules/module2/examples/Makefile

@@ -32,7 +32,7 @@ HIP_FLAGS = -std=c++17 -O2
 HIP_DEBUG_FLAGS = -std=c++17 -g
 
 # ROCm 7: Ensure hipcc can find HIP runtime by passing --rocm-path
-ROCM_PATH ?= /opt/rocm
+ROCM_PATH ?= $(shell ls -d /opt/rocm-7.0.0 2>/dev/null || ls -d /opt/rocm* 2>/dev/null | head -1 || echo /opt/rocm)
 # Auto-detect ROCm path from hipcc if headers not found
 ifeq ($(wildcard $(ROCM_PATH)/include/hip/hip_runtime.h),)
 	HIPCC_BIN := $(shell command -v hipcc 2>/dev/null)

modules/module3/examples/Makefile

@@ -32,7 +32,7 @@ HIP_FLAGS = -std=c++17 -O2
 HIP_DEBUG_FLAGS = -std=c++17 -g
 
 # ROCm 7: Ensure hipcc can find HIP runtime by passing --rocm-path
-ROCM_PATH ?= /opt/rocm
+ROCM_PATH ?= $(shell ls -d /opt/rocm-7.0.0 2>/dev/null || ls -d /opt/rocm* 2>/dev/null | head -1 || echo /opt/rocm)
 # Auto-detect ROCm path from hipcc if headers not found
 ifeq ($(wildcard $(ROCM_PATH)/include/hip/hip_runtime.h),)
 	HIPCC_BIN := $(shell command -v hipcc 2>/dev/null)

modules/module4/examples/Makefile

@@ -33,7 +33,7 @@ HIP_FLAGS = -std=c++17 -O2 -fopenmp
 HIP_DEBUG_FLAGS = -std=c++17 -g -fopenmp
 
 # ROCm 7: Ensure hipcc can find HIP runtime by passing --rocm-path
-ROCM_PATH ?= /opt/rocm
+ROCM_PATH ?= $(shell ls -d /opt/rocm-7.0.0 2>/dev/null || ls -d /opt/rocm* 2>/dev/null | head -1 || echo /opt/rocm)
 # Auto-detect ROCm path from hipcc if headers not found
 ifeq ($(wildcard $(ROCM_PATH)/include/hip/hip_runtime.h),)
 	HIPCC_BIN := $(shell command -v hipcc 2>/dev/null)

modules/module5/examples/Makefile

@@ -32,7 +32,7 @@ HIP_FLAGS = -std=c++17 -O3
 HIP_DEBUG_FLAGS = -std=c++17 -g
 
 # ROCm 7: Ensure hipcc can find HIP runtime by passing --rocm-path
-ROCM_PATH ?= /opt/rocm
+ROCM_PATH ?= $(shell ls -d /opt/rocm-7.0.0 2>/dev/null || ls -d /opt/rocm* 2>/dev/null | head -1 || echo /opt/rocm)
 # Auto-detect ROCm path from hipcc if headers not found
 ifeq ($(wildcard $(ROCM_PATH)/include/hip/hip_runtime.h),)
 	HIPCC_BIN := $(shell command -v hipcc 2>/dev/null)

modules/module6/examples/01_convolution_hip.cpp

@@ -296,7 +296,7 @@ __device__ float wavefront_reduce_sum(float val) {
 void initialize_data(float *data, int size, bool random = true) {
     for (int i = 0; i < size; i++) {
         if (random) {
-            data[i] = static_cast<float>(rand()) / RAND_MAX;
+            data[i] = static_cast<float>(rand()) / static_cast<float>(RAND_MAX);
         } else {
             data[i] = 1.0f; // Unit impulse for testing
         }

modules/module6/examples/03_histogram_hip.cpp

@@ -96,18 +96,10 @@ __global__ void histogram_wavefront_aggregation(int* input, int* histogram, int
     for (int i = idx; i < n; i += blockDim.x * gridDim.x) {
         int bin = input[i] % num_bins;
 
-        // Use ballot and popcount for wavefront aggregation
-        uint64_t mask = __ballot(1);  // All active threads in wavefront
-        int count = __popcll(mask);
-        
-        // Count how many threads in wavefront want same bin
-        uint64_t same_bin_mask = __ballot(bin == bin);  // Simplified - would need proper comparison
-        int same_bin_count = __popcll(same_bin_mask);
-        
-        // Only first thread with this bin value updates
-        if (__ffsll((unsigned long long)same_bin_mask) - 1 == lane) {
-            atomicAdd(&lds_hist[bin], same_bin_count);
-        }
+        // Simple atomic increment - removing complex wavefront aggregation for clarity
+        // In a production implementation, you would use more sophisticated wavefront
+        // aggregation by comparing bin values across the wavefront
+        atomicAdd(&lds_hist[bin], 1);
     }
 
     __syncthreads();

modules/module6/examples/05_prefix_sum_hip.cpp

@@ -182,7 +182,7 @@ __global__ void blelloch_scan_exclusive(float* input, float* output, int n) {
 // LDS bank conflict free optimization for AMD GPUs
 #define NUM_BANKS 32
 #define LOG_NUM_BANKS 5
-#define CONFLICT_FREE_OFFSET(n) (((n) >> NUM_BANKS) + ((n) >> (2 * NUM_BANKS)))
+#define CONFLICT_FREE_OFFSET(n) (((n) >> LOG_NUM_BANKS) + ((n) >> (2 * LOG_NUM_BANKS)))
 
 __global__ void blelloch_scan_lds_optimized(float* input, float* output, int n) {
     __shared__ float temp[512 + 512/NUM_BANKS];  // Extra space for conflict avoidance