Fixed the examples of module2

Author: coketaste

modules/module2/examples/01_shared_memory_transpose_hip.cpp

@@ -356,13 +356,13 @@ int main() {
     free(h_output_shared);
     free(h_output_optimized);
     free(h_output_cpu);
-    hipFree(d_input);
-    hipFree(d_output_naive);
-    hipFree(d_output_shared);
-    hipFree(d_output_optimized);
-    hipFree(d_bank_data);
-    hipEventDestroy(start);
-    hipEventDestroy(stop);
+    HIP_CHECK(hipFree(d_input));
+    HIP_CHECK(hipFree(d_output_naive));
+    HIP_CHECK(hipFree(d_output_shared));
+    HIP_CHECK(hipFree(d_output_optimized));
+    HIP_CHECK(hipFree(d_bank_data));
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(stop));
 
     printf("\nHIP shared memory transpose example completed successfully!\n");
     return 0;

modules/module2/examples/02_memory_coalescing_hip.cpp

@@ -184,9 +184,9 @@ class MemoryBenchmark {
     }
 
     ~MemoryBenchmark() {
-        hipFree(d_data);
-        hipEventDestroy(start);
-        hipEventDestroy(stop);
+        HIP_CHECK(hipFree(d_data));
+        HIP_CHECK(hipEventDestroy(start));
+        HIP_CHECK(hipEventDestroy(stop));
     }
 
     float testCoalesced(int blocks, int threads) {
@@ -382,10 +382,10 @@ void runParticleBenchmarks() {
 #endif
 
     // Cleanup
-    hipFree(soa.x); hipFree(soa.y); hipFree(soa.z);
-    hipFree(soa.vx); hipFree(soa.vy); hipFree(soa.vz);
-    hipFree(soa.mass); hipFree(aos);
-    hipEventDestroy(start); hipEventDestroy(stop);
+    HIP_CHECK(hipFree(soa.x)); HIP_CHECK(hipFree(soa.y)); HIP_CHECK(hipFree(soa.z));
+    HIP_CHECK(hipFree(soa.vx)); HIP_CHECK(hipFree(soa.vy)); HIP_CHECK(hipFree(soa.vz));
+    HIP_CHECK(hipFree(soa.mass)); HIP_CHECK(hipFree(aos));
+    HIP_CHECK(hipEventDestroy(start)); HIP_CHECK(hipEventDestroy(stop));
 }
 
 void runVectorizationBenchmark() {
@@ -439,10 +439,10 @@ void runVectorizationBenchmark() {
     printf("Vector bandwidth: %.2f GB/s\n", 
            (bytes_transferred / (1024.0 * 1024.0 * 1024.0)) / (vector_time / 1000.0));
 
-    hipFree(d_float_data);
-    hipFree(d_float4_data);
-    hipEventDestroy(start);
-    hipEventDestroy(stop);
+    HIP_CHECK(hipFree(d_float_data));
+    HIP_CHECK(hipFree(d_float4_data));
+    HIP_CHECK(hipEventDestroy(start));
+    HIP_CHECK(hipEventDestroy(stop));
 }
 
 int main() {

modules/module2/examples/03_texture_memory_hip.cpp

@@ -21,7 +21,7 @@ __global__ void textureFilterKernel(hipTextureObject_t texObj, float *output,
                 float v = (float)(y + fy + 0.5f) / height;
 
                 // Texture automatically handles boundary conditions and interpolation
-                float value = tex2D<float>(texObj, u, v);
+                float value = hipTex2D<float>(texObj, u, v);
                 sum += value;
             }
         }
@@ -42,7 +42,7 @@ __global__ void textureTranspose(hipTextureObject_t texObj, float *output,
         float v = (y + 0.5f) / height;
 
         // Fetch using texture cache
-        float value = tex2D<float>(texObj, u, v);
+        float value = hipTex2D<float>(texObj, u, v);
 
         // Write transposed
         if (y < width && x < height) {
@@ -71,7 +71,7 @@ __global__ void bilinearInterpolation(hipTextureObject_t texObj, float *output,
         float v = src_y / in_height;
 
         // Hardware bilinear interpolation
-        float interpolated = tex2D<float>(texObj, u, v);
+        float interpolated = hipTex2D<float>(texObj, u, v);
 
         output[y * out_width + x] = interpolated;
     }
@@ -134,7 +134,7 @@ __global__ void amdOptimizedTextureAccess(hipTextureObject_t texObj, float *outp
             float u = (x + 0.5f) / width;
             float v = (y + 0.5f) / height;
 
-            float value = tex2D<float>(texObj, u, v);
+            float value = hipTex2D<float>(texObj, u, v);
             output[pixel_id] = value;
         }
     }
@@ -293,10 +293,10 @@ void demonstrateTextureMemory() {
     HIP_CHECK(hipEventDestroy(start));
     HIP_CHECK(hipEventDestroy(stop));
 
-    hipFree(d_input);
-    hipFree(d_output_texture);
-    hipFree(d_output_manual);
-    hipFree(d_resized);
+    HIP_CHECK(hipFree(d_input));
+    HIP_CHECK(hipFree(d_output_texture));
+    HIP_CHECK(hipFree(d_output_manual));
+    HIP_CHECK(hipFree(d_resized));
 
     free(h_input);
     free(h_output_texture);

modules/module2/examples/04_unified_memory_hip.cpp

@@ -93,7 +93,7 @@ class UnifiedMemoryDemo {
     }
 
     ~UnifiedMemoryDemo() {
-        hipFree(data);
+        HIP_CHECK(hipFree(data));
     }
 
     void processOnGPU() {
@@ -257,7 +257,7 @@ void performanceComparison() {
 
     // Cleanup
     free(h_data);
-    hipFree(d_data);
+    HIP_CHECK(hipFree(d_data));
 }
 
 // Memory usage analysis

modules/module2/examples/05_memory_bandwidth_optimization_hip.cpp

@@ -235,8 +235,8 @@ class BandwidthTester {
     }
 
     ~BandwidthTester() {
-        hipFree(d_input);
-        hipFree(d_output);
+        HIP_CHECK(hipFree(d_input));
+        HIP_CHECK(hipFree(d_output));
     }
 
     double testBandwidth(const char* test_name, void (*kernel)(float*, float*, size_t), 
@@ -388,7 +388,7 @@ void analyzeAccessPatterns() {
         HIP_CHECK(hipEventDestroy(stop));
     }
 
-    hipFree(d_temp);
+    HIP_CHECK(hipFree(d_temp));
 }
 
 void demonstrateBandwidthOptimization() {