Skip to content

Use new accordo APIs #157

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
mawad-amd merged 7 commits into from
Nov 8, 2025
Merged

Use new accordo APIs #157

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
32c0ed5
Refactor Accordo API and integrate with clean main
mawad-amd Nov 3, 2025
096ef98
Rename AccordoValidator to Accordo and nest all classes
mawad-amd Nov 3, 2025
e611379
Apply Ruff auto-fixes
github-actions[bot] Nov 4, 2025
f09c2fd
Move history before reporting results
mawad-amd Nov 8, 2025
16ec3fe
Apply Ruff auto-fixes
github-actions[bot] Nov 8, 2025
10069ef
Add licence
mawad-amd Nov 8, 2025
13d424f
Fix missing args
mawad-amd Nov 8, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions examples/access_pattern/uncoalesced/uncoalesced.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ __global__ void matrix_transpose(const T* __restrict__ in,
}
int main() {
using T = __hip_bfloat16;

const int width = 1024;
const int height = 1024;
const int size = width * height;
Expand Down Expand Up @@ -97,7 +97,7 @@ int main() {

std::cout << (correct ? "Transpose correct ✅" : "Transpose incorrect ❌") << "\n";

hipFree(d_in);
hipFree(d_out);
hip_try(hipFree(d_in));
hip_try(hipFree(d_out));
return 0;
}
22 changes: 16 additions & 6 deletions examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,16 @@ SOFTWARE.
#include <hip/hip_runtime.h>
#include <vector>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define TILE_DIM 16

__global__ void matrixTransposeShared_0(float* out,
Expand Down Expand Up @@ -106,9 +116,9 @@ void runTranspose(int width, int height) {

float* d_in;
float* d_out;
hipMalloc(&d_in, width * height * sizeof(float));
hipMalloc(&d_out, width * height * sizeof(float));
hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice);
hip_try(hipMalloc(&d_in, width * height * sizeof(float)));
hip_try(hipMalloc(&d_out, width * height * sizeof(float)));
hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice));
dim3 blockSize(TILE_DIM, TILE_DIM);
dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM);

Expand All @@ -120,10 +130,10 @@ void runTranspose(int width, int height) {
if (status != hipSuccess) {
std::terminate();
}
hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost);
hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost));

hipFree(d_in);
hipFree(d_out);
hip_try(hipFree(d_in));
hip_try(hipFree(d_out));
}

int main() {
Expand Down
22 changes: 16 additions & 6 deletions examples/bank_conflict/matrix_transpose/matrix_transpose.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,16 @@ SOFTWARE.
#include <cstdlib>
#include <iostream>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define TILE_DIM 16

__global__ void matrixTransposeShared(float* out,
Expand Down Expand Up @@ -62,9 +72,9 @@ void runTranspose(int width, int height) {

float* d_in;
float* d_out;
hipMalloc(&d_in, width * height * sizeof(float));
hipMalloc(&d_out, width * height * sizeof(float));
hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice);
hip_try(hipMalloc(&d_in, width * height * sizeof(float)));
hip_try(hipMalloc(&d_out, width * height * sizeof(float)));
hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice));
dim3 blockSize(TILE_DIM, TILE_DIM);
dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM);

Expand All @@ -74,10 +84,10 @@ void runTranspose(int width, int height) {
if(status != hipSuccess){
std::terminate();
}
hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost);
hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost));

hipFree(d_in);
hipFree(d_out);
hip_try(hipFree(d_in));
hip_try(hipFree(d_out));
}

int main(int argc, char* argv[]) {
Expand Down
24 changes: 17 additions & 7 deletions examples/bank_conflict/reduce/reduce.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,16 @@ SOFTWARE.
#include <cstdio>
#include <hip/hip_runtime.h>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define BLOCK_SIZE 256

__global__ void reduce_kernel(const float *d_in, float *d_out, int n) {
Expand Down Expand Up @@ -61,22 +71,22 @@ int main() {
}

float *d_in = nullptr, *d_out = nullptr;
hipMalloc((void **)&d_in, size);
hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE));
hip_try(hipMalloc((void **)&d_in, size));
hip_try(hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE)));

hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice);
hip_try(hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice));

int gridSize = (numElements + BLOCK_SIZE - 1) / BLOCK_SIZE;
reduce_kernel<<<gridSize, BLOCK_SIZE>>>(d_in, d_out, numElements);
hipDeviceSynchronize();
hip_try(hipDeviceSynchronize());

hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost);
hip_try(hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost));

printf("First block sum: %f\n", h_out[0]);

// Free resources.
hipFree(d_in);
hipFree(d_out);
hip_try(hipFree(d_in));
hip_try(hipFree(d_out));
free(h_in);
free(h_out);

Expand Down
16 changes: 13 additions & 3 deletions examples/bank_conflict/synthetic/synthetic.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -24,6 +24,16 @@ SOFTWARE.

#include <hip/hip_runtime.h>
#include <iostream>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define BLOCK_SIZE 256
#define LDS_SIZE 256
Expand Down Expand Up @@ -52,20 +62,20 @@ int main() {

// Allocate memory on the device
float* d_out;
hipMalloc(&d_out, BLOCK_SIZE * sizeof(float));
hip_try(hipMalloc(&d_out, BLOCK_SIZE * sizeof(float)));

dim3 blockSize(BLOCK_SIZE);
dim3 gridSize(1);
hipLaunchKernelGGL(bankConflictKernel, gridSize, blockSize, 0, 0, d_out);

// Copy the result back to the host
hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost);
hip_try(hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost));

for (int i = 0; i < BLOCK_SIZE; ++i) {
std::cout << "h_out[" << i << "] = " << h_out[i] << std::endl;
}

hipFree(d_out);
hip_try(hipFree(d_out));

return 0;
}
24 changes: 17 additions & 7 deletions examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,16 @@ SOFTWARE.
#include <numeric>
#include <vector>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define TILE_DIM 16

__global__ void matrixTransposeShared(float* out,
Expand Down Expand Up @@ -75,9 +85,9 @@ void runKernels(int width, int height) {
std::iota(h_data.begin(), h_data.end(), 0.0f);

float *d_data, *d_out;
hipMalloc(&d_data, num_elements * sizeof(float));
hipMalloc(&d_out, num_elements * sizeof(float));
hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice);
hip_try(hipMalloc(&d_data, num_elements * sizeof(float)));
hip_try(hipMalloc(&d_out, num_elements * sizeof(float)));
hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice));

// --- 1. Scale all elements by 2.0
int blockSize = 256;
Expand All @@ -93,11 +103,11 @@ void runKernels(int width, int height) {
matrixTransposeShared<<<gridDim, blockDim>>>(d_out, d_data, width, height);

// --- Copy result back
hipDeviceSynchronize();
hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost);
hip_try(hipDeviceSynchronize());
hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost));

hipFree(d_data);
hipFree(d_out);
hip_try(hipFree(d_data));
hip_try(hipFree(d_out));

// Print a few values
std::cout << "Result (first 10 elements):\n";
Expand Down
24 changes: 17 additions & 7 deletions examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,16 @@ SOFTWARE.
#include <numeric>
#include <vector>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


#define TILE_DIM 16

template <typename T>
Expand Down Expand Up @@ -76,9 +86,9 @@ void runKernels(int width, int height) {
std::iota(h_data.begin(), h_data.end(), static_cast<T>(0));

T *d_data, *d_out;
hipMalloc(&d_data, num_elements * sizeof(T));
hipMalloc(&d_out, num_elements * sizeof(T));
hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice);
hip_try(hipMalloc(&d_data, num_elements * sizeof(T)));
hip_try(hipMalloc(&d_out, num_elements * sizeof(T)));
hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice));

// --- 1. Scale all elements
int blockSize = 256;
Expand All @@ -93,11 +103,11 @@ void runKernels(int width, int height) {
dim3 gridDim((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM);
matrixTransposeShared<T><<<gridDim, blockDim>>>(d_out, d_data, width, height);

hipDeviceSynchronize();
hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost);
hip_try(hipDeviceSynchronize());
hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost));

hipFree(d_data);
hipFree(d_out);
hip_try(hipFree(d_data));
hip_try(hipFree(d_out));

std::cout << "Result (first 10 elements):\n";
for (int i = 0; i < 10; ++i)
Expand Down
44 changes: 27 additions & 17 deletions examples/basic/vector_add/vector_add.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,16 @@ SOFTWARE.
#include <hip/hip_runtime.h>
#include <iostream>

#define hip_try(expr) \
do { \
hipError_t err = (expr); \
if (err != hipSuccess) { \
const char* msg = hipGetErrorString(err); \
throw std::runtime_error(std::string("HIP error: ") + msg); \
} \
} while (0)


__global__ void vector_add(const float* a, const float* b, float* c, size_t n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n)
Expand All @@ -46,38 +56,38 @@ int main() {
}

float *d_a, *d_b, *d_c;
hipMalloc(&d_a, size);
hipMalloc(&d_b, size);
hipMalloc(&d_c, size);
hip_try(hipMalloc(&d_a, size));
hip_try(hipMalloc(&d_b, size));
hip_try(hipMalloc(&d_c, size));

hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice);
hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice);
hip_try(hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice));
hip_try(hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice));

const int threadsPerBlock = 256;
const int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;

// Create HIP events
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
hip_try(hipEventCreate(&start));
hip_try(hipEventCreate(&stop));

// Record start
hipEventRecord(start, 0);
hip_try(hipEventRecord(start, 0));

// Launch kernel
hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threadsPerBlock), 0, 0, d_a, d_b, d_c, N);

// Record stop
hipEventRecord(stop, 0);
hipEventSynchronize(stop);
hip_try(hipEventRecord(stop, 0));
hip_try(hipEventSynchronize(stop));

// Calculate elapsed time
float milliseconds = 0;
hipEventElapsedTime(&milliseconds, start, stop);
hip_try(hipEventElapsedTime(&milliseconds, start, stop));
std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl;

// Copy result back to host
hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost);
hip_try(hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost));

// Print some results
for (size_t i = 0; i < 5; ++i)
Expand All @@ -87,11 +97,11 @@ int main() {
delete[] h_a;
delete[] h_b;
delete[] h_c;
hipFree(d_a);
hipFree(d_b);
hipFree(d_c);
hipEventDestroy(start);
hipEventDestroy(stop);
hip_try(hipFree(d_a));
hip_try(hipFree(d_b));
hip_try(hipFree(d_c));
hip_try(hipEventDestroy(start));
hip_try(hipEventDestroy(stop));

return 0;
}
Loading
Loading